transform: Don't unroll arrays in DecomposeMemoryAccess
Arrays can be extremely large, and having the load and store functions unroll the elements can make the complier explode.
Fixed: chromium:1229233
Change-Id: Ieb5654254e16f5ce724a205d21d954ef9a0cd053
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/58382
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@chromium.org>
Reviewed-by: David Neto <dneto@google.com>
Auto-Submit: Ben Clayton <bclayton@google.com>
diff --git a/src/transform/decompose_memory_access.cc b/src/transform/decompose_memory_access.cc
index 06251de..f06b3ae 100644
--- a/src/transform/decompose_memory_access.cc
+++ b/src/transform/decompose_memory_access.cc
@@ -303,6 +303,10 @@
/// State holds the current transform state
struct DecomposeMemoryAccess::State {
+ /// The clone context
+ CloneContext& ctx;
+ /// Alias to `*ctx.dst`
+ ProgramBuilder& b;
/// Map of AST expression to storage or uniform buffer access
/// This map has entries added when encountered, and removed when outer
/// expressions chain the access.
@@ -322,6 +326,10 @@
/// Allocations for offsets
BlockAllocator<Offset> offsets_;
+ /// Constructor
+ /// @param context the CloneContext
+ explicit State(CloneContext& context) : ctx(context), b(*ctx.dst) {}
+
/// @param offset the offset value to wrap in an Offset
/// @returns an Offset for the given literal value
const Offset* ToOffset(uint32_t offset) {
@@ -440,13 +448,11 @@
/// of type `el_ty` from a storage or uniform buffer of type `buf_ty`.
/// The emitted function has the signature:
/// `fn load(buf : buf_ty, offset : u32) -> el_ty`
- /// @param ctx the CloneContext
/// @param buf_ty the storage or uniform buffer type
/// @param el_ty the storage or uniform buffer element type
/// @param var_user the variable user
/// @return the name of the function that performs the load
- Symbol LoadFunc(CloneContext& ctx,
- const sem::Type* buf_ty,
+ Symbol LoadFunc(const sem::Type* buf_ty,
const sem::Type* el_ty,
const sem::VariableUser* var_user) {
auto storage_class = var_user->Variable()->StorageClass();
@@ -454,70 +460,89 @@
load_funcs, LoadStoreKey{storage_class, buf_ty, el_ty}, [&] {
auto* buf_ast_ty = CreateASTTypeFor(ctx, buf_ty);
auto* disable_validation =
- ctx.dst->ASTNodes().Create<ast::DisableValidationDecoration>(
- ctx.dst->ID(), ast::DisabledValidation::
- kIgnoreConstructibleFunctionParameter);
+ b.ASTNodes().Create<ast::DisableValidationDecoration>(
+ b.ID(), ast::DisabledValidation::
+ kIgnoreConstructibleFunctionParameter);
ast::VariableList params = {
// Note: The buffer parameter requires the StorageClass in
// order for HLSL to emit this as a ByteAddressBuffer or cbuffer
// array.
- ctx.dst->create<ast::Variable>(
- ctx.dst->Sym("buffer"), storage_class,
- var_user->Variable()->Access(), buf_ast_ty, true, nullptr,
- ast::DecorationList{disable_validation}),
- ctx.dst->Param("offset", ctx.dst->ty.u32()),
+ b.create<ast::Variable>(b.Sym("buffer"), storage_class,
+ var_user->Variable()->Access(),
+ buf_ast_ty, true, nullptr,
+ ast::DecorationList{disable_validation}),
+ b.Param("offset", b.ty.u32()),
};
- ast::Function* func = nullptr;
+ auto name = b.Sym();
+
if (auto* intrinsic =
IntrinsicLoadFor(ctx.dst, storage_class, el_ty)) {
auto* el_ast_ty = CreateASTTypeFor(ctx, el_ty);
- func = ctx.dst->create<ast::Function>(
- ctx.dst->Sym(), params, el_ast_ty, nullptr,
+ auto* func = b.create<ast::Function>(
+ name, params, el_ast_ty, nullptr,
ast::DecorationList{
intrinsic,
- ctx.dst->ASTNodes()
- .Create<ast::DisableValidationDecoration>(
- ctx.dst->ID(),
- ast::DisabledValidation::kFunctionHasNoBody),
+ b.ASTNodes().Create<ast::DisableValidationDecoration>(
+ b.ID(), ast::DisabledValidation::kFunctionHasNoBody),
},
ast::DecorationList{});
+ b.AST().AddFunction(func);
+ } else if (auto* arr_ty = el_ty->As<sem::Array>()) {
+ // fn load_func(buf : buf_ty, offset : u32) -> array<T, N> {
+ // var arr : array<T, N>;
+ // for (var i = 0u; i < array_count; i = i + 1) {
+ // arr[i] = el_load_func(buf, offset + i * array_stride)
+ // }
+ // return arr;
+ // }
+ auto load =
+ LoadFunc(buf_ty, arr_ty->ElemType()->UnwrapRef(), var_user);
+ auto* arr =
+ b.Var(b.Symbols().New("arr"), CreateASTTypeFor(ctx, arr_ty));
+ auto* i = b.Var(b.Symbols().New("i"), nullptr, b.Expr(0u));
+ auto* for_init = b.Decl(i);
+ auto* for_cond = b.create<ast::BinaryExpression>(
+ ast::BinaryOp::kLessThan, b.Expr(i), b.Expr(arr_ty->Count()));
+ auto* for_cont = b.Assign(i, b.Add(i, 1u));
+ auto* arr_el = b.IndexAccessor(arr, i);
+ auto* el_offset =
+ b.Add(b.Expr("offset"), b.Mul(i, arr_ty->Stride()));
+ auto* el_val = b.Call(load, "buffer", el_offset);
+ auto* for_loop = b.For(for_init, for_cond, for_cont,
+ b.Block(b.Assign(arr_el, el_val)));
+
+ b.Func(name, params, CreateASTTypeFor(ctx, arr_ty),
+ {
+ b.Decl(arr),
+ for_loop,
+ b.Return(arr),
+ });
} else {
ast::ExpressionList values;
if (auto* mat_ty = el_ty->As<sem::Matrix>()) {
auto* vec_ty = mat_ty->ColumnType();
- Symbol load = LoadFunc(ctx, buf_ty, vec_ty, var_user);
+ Symbol load = LoadFunc(buf_ty, vec_ty, var_user);
for (uint32_t i = 0; i < mat_ty->columns(); i++) {
- auto* offset =
- ctx.dst->Add("offset", i * MatrixColumnStride(mat_ty));
- values.emplace_back(ctx.dst->Call(load, "buffer", offset));
+ auto* offset = b.Add("offset", i * MatrixColumnStride(mat_ty));
+ values.emplace_back(b.Call(load, "buffer", offset));
}
} else if (auto* str = el_ty->As<sem::Struct>()) {
for (auto* member : str->Members()) {
- auto* offset = ctx.dst->Add("offset", member->Offset());
- Symbol load = LoadFunc(ctx, buf_ty, member->Type()->UnwrapRef(),
- var_user);
- values.emplace_back(ctx.dst->Call(load, "buffer", offset));
- }
- } else if (auto* arr = el_ty->As<sem::Array>()) {
- for (uint32_t i = 0; i < arr->Count(); i++) {
- auto* offset = ctx.dst->Add("offset", arr->Stride() * i);
- Symbol load = LoadFunc(ctx, buf_ty,
- arr->ElemType()->UnwrapRef(), var_user);
- values.emplace_back(ctx.dst->Call(load, "buffer", offset));
+ auto* offset = b.Add("offset", member->Offset());
+ Symbol load =
+ LoadFunc(buf_ty, member->Type()->UnwrapRef(), var_user);
+ values.emplace_back(b.Call(load, "buffer", offset));
}
}
- auto* el_ast_ty = CreateASTTypeFor(ctx, el_ty);
- func = ctx.dst->create<ast::Function>(
- ctx.dst->Sym(), params, el_ast_ty,
- ctx.dst->Block(ctx.dst->Return(
- ctx.dst->create<ast::TypeConstructorExpression>(
- CreateASTTypeFor(ctx, el_ty), values))),
- ast::DecorationList{}, ast::DecorationList{});
+ b.Func(name, params, CreateASTTypeFor(ctx, el_ty),
+ {
+ b.Return(b.create<ast::TypeConstructorExpression>(
+ CreateASTTypeFor(ctx, el_ty), values)),
+ });
}
- ctx.dst->AST().AddFunction(func);
- return func->symbol();
+ return name;
});
}
@@ -525,13 +550,11 @@
/// element of type `el_ty` to a storage buffer of type `buf_ty`.
/// The function has the signature:
/// `fn store(buf : buf_ty, offset : u32, value : el_ty)`
- /// @param ctx the CloneContext
/// @param buf_ty the storage buffer type
/// @param el_ty the storage buffer element type
/// @param var_user the variable user
/// @return the name of the function that performs the store
- Symbol StoreFunc(CloneContext& ctx,
- const sem::Type* buf_ty,
+ Symbol StoreFunc(const sem::Type* buf_ty,
const sem::Type* el_ty,
const sem::VariableUser* var_user) {
auto storage_class = var_user->Variable()->StorageClass();
@@ -540,75 +563,87 @@
auto* buf_ast_ty = CreateASTTypeFor(ctx, buf_ty);
auto* el_ast_ty = CreateASTTypeFor(ctx, el_ty);
auto* disable_validation =
- ctx.dst->ASTNodes().Create<ast::DisableValidationDecoration>(
- ctx.dst->ID(), ast::DisabledValidation::
- kIgnoreConstructibleFunctionParameter);
+ b.ASTNodes().Create<ast::DisableValidationDecoration>(
+ b.ID(), ast::DisabledValidation::
+ kIgnoreConstructibleFunctionParameter);
ast::VariableList params{
// Note: The buffer parameter requires the StorageClass in
// order for HLSL to emit this as a ByteAddressBuffer.
- ctx.dst->create<ast::Variable>(
- ctx.dst->Sym("buffer"), storage_class,
- var_user->Variable()->Access(), buf_ast_ty, true, nullptr,
- ast::DecorationList{disable_validation}),
- ctx.dst->Param("offset", ctx.dst->ty.u32()),
- ctx.dst->Param("value", el_ast_ty),
+ b.create<ast::Variable>(b.Sym("buffer"), storage_class,
+ var_user->Variable()->Access(),
+ buf_ast_ty, true, nullptr,
+ ast::DecorationList{disable_validation}),
+ b.Param("offset", b.ty.u32()),
+ b.Param("value", el_ast_ty),
};
- ast::Function* func = nullptr;
+
+ auto name = b.Sym();
+
if (auto* intrinsic =
IntrinsicStoreFor(ctx.dst, storage_class, el_ty)) {
- func = ctx.dst->create<ast::Function>(
- ctx.dst->Sym(), params, ctx.dst->ty.void_(), nullptr,
+ auto* func = b.create<ast::Function>(
+ name, params, b.ty.void_(), nullptr,
ast::DecorationList{
intrinsic,
- ctx.dst->ASTNodes()
- .Create<ast::DisableValidationDecoration>(
- ctx.dst->ID(),
- ast::DisabledValidation::kFunctionHasNoBody),
+ b.ASTNodes().Create<ast::DisableValidationDecoration>(
+ b.ID(), ast::DisabledValidation::kFunctionHasNoBody),
},
ast::DecorationList{});
-
+ b.AST().AddFunction(func);
} else {
ast::StatementList body;
- if (auto* mat_ty = el_ty->As<sem::Matrix>()) {
+ if (auto* arr_ty = el_ty->As<sem::Array>()) {
+ // fn store_func(buf : buf_ty, offset : u32, value : el_ty) {
+ // var array = value; // No dynamic indexing on constant arrays
+ // for (var i = 0u; i < array_count; i = i + 1) {
+ // arr[i] = el_store_func(buf, offset + i * array_stride,
+ // value[i])
+ // }
+ // return arr;
+ // }
+ auto* array =
+ b.Var(b.Symbols().New("array"), nullptr, b.Expr("value"));
+ auto store =
+ StoreFunc(buf_ty, arr_ty->ElemType()->UnwrapRef(), var_user);
+ auto* i = b.Var(b.Symbols().New("i"), nullptr, b.Expr(0u));
+ auto* for_init = b.Decl(i);
+ auto* for_cond = b.create<ast::BinaryExpression>(
+ ast::BinaryOp::kLessThan, b.Expr(i), b.Expr(arr_ty->Count()));
+ auto* for_cont = b.Assign(i, b.Add(i, 1u));
+ auto* arr_el = b.IndexAccessor(array, i);
+ auto* el_offset =
+ b.Add(b.Expr("offset"), b.Mul(i, arr_ty->Stride()));
+ auto* store_stmt = b.create<ast::CallStatement>(
+ b.Call(store, "buffer", el_offset, arr_el));
+ auto* for_loop =
+ b.For(for_init, for_cond, for_cont, b.Block(store_stmt));
+
+ body = {b.Decl(array), for_loop};
+ } else if (auto* mat_ty = el_ty->As<sem::Matrix>()) {
auto* vec_ty = mat_ty->ColumnType();
- Symbol store = StoreFunc(ctx, buf_ty, vec_ty, var_user);
+ Symbol store = StoreFunc(buf_ty, vec_ty, var_user);
for (uint32_t i = 0; i < mat_ty->columns(); i++) {
- auto* offset =
- ctx.dst->Add("offset", i * MatrixColumnStride(mat_ty));
- auto* access = ctx.dst->IndexAccessor("value", i);
- auto* call = ctx.dst->Call(store, "buffer", offset, access);
- body.emplace_back(ctx.dst->create<ast::CallStatement>(call));
+ auto* offset = b.Add("offset", i * MatrixColumnStride(mat_ty));
+ auto* access = b.IndexAccessor("value", i);
+ auto* call = b.Call(store, "buffer", offset, access);
+ body.emplace_back(b.create<ast::CallStatement>(call));
}
} else if (auto* str = el_ty->As<sem::Struct>()) {
for (auto* member : str->Members()) {
- auto* offset = ctx.dst->Add("offset", member->Offset());
- auto* access = ctx.dst->MemberAccessor(
+ auto* offset = b.Add("offset", member->Offset());
+ auto* access = b.MemberAccessor(
"value", ctx.Clone(member->Declaration()->symbol()));
- Symbol store = StoreFunc(ctx, buf_ty,
- member->Type()->UnwrapRef(), var_user);
- auto* call = ctx.dst->Call(store, "buffer", offset, access);
- body.emplace_back(ctx.dst->create<ast::CallStatement>(call));
- }
- } else if (auto* arr = el_ty->As<sem::Array>()) {
- for (uint32_t i = 0; i < arr->Count(); i++) {
- auto* offset = ctx.dst->Add("offset", arr->Stride() * i);
- auto* access =
- ctx.dst->IndexAccessor("value", ctx.dst->Expr(i));
- Symbol store = StoreFunc(
- ctx, buf_ty, arr->ElemType()->UnwrapRef(), var_user);
- auto* call = ctx.dst->Call(store, "buffer", offset, access);
- body.emplace_back(ctx.dst->create<ast::CallStatement>(call));
+ Symbol store =
+ StoreFunc(buf_ty, member->Type()->UnwrapRef(), var_user);
+ auto* call = b.Call(store, "buffer", offset, access);
+ body.emplace_back(b.create<ast::CallStatement>(call));
}
}
- func = ctx.dst->create<ast::Function>(
- ctx.dst->Sym(), params, ctx.dst->ty.void_(),
- ctx.dst->Block(body), ast::DecorationList{},
- ast::DecorationList{});
+ b.Func(name, params, b.ty.void_(), body);
}
- ctx.dst->AST().AddFunction(func);
- return func->symbol();
+ return name;
});
}
@@ -616,14 +651,12 @@
/// atomic operation from a storage buffer of type `buf_ty`. The function has
/// the signature:
// `fn atomic_op(buf : buf_ty, offset : u32, ...) -> T`
- /// @param ctx the CloneContext
/// @param buf_ty the storage buffer type
/// @param el_ty the storage buffer element type
/// @param intrinsic the atomic intrinsic
/// @param var_user the variable user
/// @return the name of the function that performs the load
- Symbol AtomicFunc(CloneContext& ctx,
- const sem::Type* buf_ty,
+ Symbol AtomicFunc(const sem::Type* buf_ty,
const sem::Type* el_ty,
const sem::Intrinsic* intrinsic,
const sem::VariableUser* var_user) {
@@ -631,8 +664,8 @@
return utils::GetOrCreate(atomic_funcs, AtomicKey{buf_ty, el_ty, op}, [&] {
auto* buf_ast_ty = CreateASTTypeFor(ctx, buf_ty);
auto* disable_validation =
- ctx.dst->ASTNodes().Create<ast::DisableValidationDecoration>(
- ctx.dst->ID(),
+ b.ASTNodes().Create<ast::DisableValidationDecoration>(
+ b.ID(),
ast::DisabledValidation::kIgnoreConstructibleFunctionParameter);
// The first parameter to all WGSL atomics is the expression to the
// atomic. This is replaced with two parameters: the buffer and offset.
@@ -640,38 +673,38 @@
ast::VariableList params = {
// Note: The buffer parameter requires the kStorage StorageClass in
// order for HLSL to emit this as a ByteAddressBuffer.
- ctx.dst->create<ast::Variable>(
- ctx.dst->Sym("buffer"), ast::StorageClass::kStorage,
- var_user->Variable()->Access(), buf_ast_ty, true, nullptr,
- ast::DecorationList{disable_validation}),
- ctx.dst->Param("offset", ctx.dst->ty.u32()),
+ b.create<ast::Variable>(b.Sym("buffer"), ast::StorageClass::kStorage,
+ var_user->Variable()->Access(), buf_ast_ty,
+ true, nullptr,
+ ast::DecorationList{disable_validation}),
+ b.Param("offset", b.ty.u32()),
};
// Other parameters are copied as-is:
for (size_t i = 1; i < intrinsic->Parameters().size(); i++) {
auto& param = intrinsic->Parameters()[i];
auto* ty = CreateASTTypeFor(ctx, param.type);
- params.emplace_back(ctx.dst->Param("param_" + std::to_string(i), ty));
+ params.emplace_back(b.Param("param_" + std::to_string(i), ty));
}
auto* atomic = IntrinsicAtomicFor(ctx.dst, op, el_ty);
if (atomic == nullptr) {
- TINT_ICE(Transform, ctx.dst->Diagnostics())
+ TINT_ICE(Transform, b.Diagnostics())
<< "IntrinsicAtomicFor() returned nullptr for op " << op
<< " and type " << el_ty->type_name();
}
auto* ret_ty = CreateASTTypeFor(ctx, intrinsic->ReturnType());
- auto* func = ctx.dst->create<ast::Function>(
- ctx.dst->Sym(), params, ret_ty, nullptr,
+ auto* func = b.create<ast::Function>(
+ b.Sym(), params, ret_ty, nullptr,
ast::DecorationList{
atomic,
- ctx.dst->ASTNodes().Create<ast::DisableValidationDecoration>(
- ctx.dst->ID(), ast::DisabledValidation::kFunctionHasNoBody),
+ b.ASTNodes().Create<ast::DisableValidationDecoration>(
+ b.ID(), ast::DisabledValidation::kFunctionHasNoBody),
},
ast::DecorationList{});
- ctx.dst->AST().AddFunction(func);
+ b.AST().AddFunction(func);
return func->symbol();
});
}
@@ -777,7 +810,7 @@
void DecomposeMemoryAccess::Run(CloneContext& ctx, const DataMap&, DataMap&) {
auto& sem = ctx.src->Sem();
- State state;
+ State state(ctx);
// Scan the AST nodes for storage and uniform buffer accesses. Complex
// expression chains (e.g. `storage_buffer.foo.bar[20].x`) are handled by
@@ -908,7 +941,7 @@
auto* buf_ty = access.var->Type()->UnwrapRef();
auto* el_ty = access.type->UnwrapRef()->As<sem::Atomic>()->Type();
Symbol func =
- state.AtomicFunc(ctx, buf_ty, el_ty, intrinsic,
+ state.AtomicFunc(buf_ty, el_ty, intrinsic,
access.var->As<sem::VariableUser>());
ast::ExpressionList args{ctx.Clone(buf), offset};
@@ -937,8 +970,8 @@
auto* offset = access.offset->Build(ctx);
auto* buf_ty = access.var->Type()->UnwrapRef();
auto* el_ty = access.type->UnwrapRef();
- Symbol func = state.LoadFunc(ctx, buf_ty, el_ty,
- access.var->As<sem::VariableUser>());
+ Symbol func =
+ state.LoadFunc(buf_ty, el_ty, access.var->As<sem::VariableUser>());
return ctx.dst->Call(func, ctx.CloneWithoutTransform(buf), offset);
});
}
@@ -951,7 +984,7 @@
auto* buf_ty = store.target.var->Type()->UnwrapRef();
auto* el_ty = store.target.type->UnwrapRef();
auto* value = store.assignment->rhs();
- Symbol func = state.StoreFunc(ctx, buf_ty, el_ty,
+ Symbol func = state.StoreFunc(buf_ty, el_ty,
store.target.var->As<sem::VariableUser>());
auto* call = ctx.dst->Call(func, ctx.CloneWithoutTransform(buf), offset,
ctx.Clone(value));
diff --git a/src/transform/decompose_memory_access_test.cc b/src/transform/decompose_memory_access_test.cc
index 5e9e60d..58466c4 100644
--- a/src/transform/decompose_memory_access_test.cc
+++ b/src/transform/decompose_memory_access_test.cc
@@ -181,7 +181,11 @@
}
fn tint_symbol_21([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> array<vec3<f32>, 2> {
- return array<vec3<f32>, 2>(tint_symbol_8(buffer, (offset + 0u)), tint_symbol_8(buffer, (offset + 16u)));
+ var arr : array<vec3<f32>, 2>;
+ for(var i_1 = 0u; (i_1 < 2u); i_1 = (i_1 + 1u)) {
+ arr[i_1] = tint_symbol_8(buffer, (offset + (i_1 * 16u)));
+ }
+ return arr;
}
[[stage(compute), workgroup_size(1)]]
@@ -375,7 +379,11 @@
}
fn tint_symbol_21([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : UB, offset : u32) -> array<vec3<f32>, 2> {
- return array<vec3<f32>, 2>(tint_symbol_8(buffer, (offset + 0u)), tint_symbol_8(buffer, (offset + 16u)));
+ var arr : array<vec3<f32>, 2>;
+ for(var i_1 = 0u; (i_1 < 2u); i_1 = (i_1 + 1u)) {
+ arr[i_1] = tint_symbol_8(buffer, (offset + (i_1 * 16u)));
+ }
+ return arr;
}
[[stage(compute), workgroup_size(1)]]
@@ -587,8 +595,10 @@
}
fn tint_symbol_21([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : array<vec3<f32>, 2>) {
- tint_symbol_8(buffer, (offset + 0u), value[0u]);
- tint_symbol_8(buffer, (offset + 16u), value[1u]);
+ var array = value;
+ for(var i_1 = 0u; (i_1 < 2u); i_1 = (i_1 + 1u)) {
+ tint_symbol_8(buffer, (offset + (i_1 * 16u)), array[i_1]);
+ }
}
[[stage(compute), workgroup_size(1)]]
@@ -689,88 +699,92 @@
[[group(0), binding(0)]] var<storage, read_write> sb : SB;
[[internal(intrinsic_load_storage_i32), internal(disable_validation__function_has_no_body)]]
-fn tint_symbol([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> i32
+fn tint_symbol_1([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> i32
[[internal(intrinsic_load_storage_u32), internal(disable_validation__function_has_no_body)]]
-fn tint_symbol_1([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> u32
+fn tint_symbol_2([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> u32
[[internal(intrinsic_load_storage_f32), internal(disable_validation__function_has_no_body)]]
-fn tint_symbol_2([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> f32
+fn tint_symbol_3([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> f32
[[internal(intrinsic_load_storage_vec2_i32), internal(disable_validation__function_has_no_body)]]
-fn tint_symbol_3([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec2<i32>
+fn tint_symbol_4([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec2<i32>
[[internal(intrinsic_load_storage_vec2_u32), internal(disable_validation__function_has_no_body)]]
-fn tint_symbol_4([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec2<u32>
+fn tint_symbol_5([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec2<u32>
[[internal(intrinsic_load_storage_vec2_f32), internal(disable_validation__function_has_no_body)]]
-fn tint_symbol_5([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec2<f32>
+fn tint_symbol_6([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec2<f32>
[[internal(intrinsic_load_storage_vec3_i32), internal(disable_validation__function_has_no_body)]]
-fn tint_symbol_6([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec3<i32>
+fn tint_symbol_7([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec3<i32>
[[internal(intrinsic_load_storage_vec3_u32), internal(disable_validation__function_has_no_body)]]
-fn tint_symbol_7([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec3<u32>
+fn tint_symbol_8([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec3<u32>
[[internal(intrinsic_load_storage_vec3_f32), internal(disable_validation__function_has_no_body)]]
-fn tint_symbol_8([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec3<f32>
+fn tint_symbol_9([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec3<f32>
[[internal(intrinsic_load_storage_vec4_i32), internal(disable_validation__function_has_no_body)]]
-fn tint_symbol_9([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec4<i32>
+fn tint_symbol_10([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec4<i32>
[[internal(intrinsic_load_storage_vec4_u32), internal(disable_validation__function_has_no_body)]]
-fn tint_symbol_10([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec4<u32>
+fn tint_symbol_11([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec4<u32>
[[internal(intrinsic_load_storage_vec4_f32), internal(disable_validation__function_has_no_body)]]
-fn tint_symbol_11([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec4<f32>
+fn tint_symbol_12([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec4<f32>
-fn tint_symbol_12([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat2x2<f32> {
- return mat2x2<f32>(tint_symbol_5(buffer, (offset + 0u)), tint_symbol_5(buffer, (offset + 8u)));
+fn tint_symbol_13([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat2x2<f32> {
+ return mat2x2<f32>(tint_symbol_6(buffer, (offset + 0u)), tint_symbol_6(buffer, (offset + 8u)));
}
-fn tint_symbol_13([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat2x3<f32> {
- return mat2x3<f32>(tint_symbol_8(buffer, (offset + 0u)), tint_symbol_8(buffer, (offset + 16u)));
+fn tint_symbol_14([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat2x3<f32> {
+ return mat2x3<f32>(tint_symbol_9(buffer, (offset + 0u)), tint_symbol_9(buffer, (offset + 16u)));
}
-fn tint_symbol_14([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat2x4<f32> {
- return mat2x4<f32>(tint_symbol_11(buffer, (offset + 0u)), tint_symbol_11(buffer, (offset + 16u)));
+fn tint_symbol_15([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat2x4<f32> {
+ return mat2x4<f32>(tint_symbol_12(buffer, (offset + 0u)), tint_symbol_12(buffer, (offset + 16u)));
}
-fn tint_symbol_15([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat3x2<f32> {
- return mat3x2<f32>(tint_symbol_5(buffer, (offset + 0u)), tint_symbol_5(buffer, (offset + 8u)), tint_symbol_5(buffer, (offset + 16u)));
+fn tint_symbol_16([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat3x2<f32> {
+ return mat3x2<f32>(tint_symbol_6(buffer, (offset + 0u)), tint_symbol_6(buffer, (offset + 8u)), tint_symbol_6(buffer, (offset + 16u)));
}
-fn tint_symbol_16([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat3x3<f32> {
- return mat3x3<f32>(tint_symbol_8(buffer, (offset + 0u)), tint_symbol_8(buffer, (offset + 16u)), tint_symbol_8(buffer, (offset + 32u)));
+fn tint_symbol_17([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat3x3<f32> {
+ return mat3x3<f32>(tint_symbol_9(buffer, (offset + 0u)), tint_symbol_9(buffer, (offset + 16u)), tint_symbol_9(buffer, (offset + 32u)));
}
-fn tint_symbol_17([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat3x4<f32> {
- return mat3x4<f32>(tint_symbol_11(buffer, (offset + 0u)), tint_symbol_11(buffer, (offset + 16u)), tint_symbol_11(buffer, (offset + 32u)));
+fn tint_symbol_18([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat3x4<f32> {
+ return mat3x4<f32>(tint_symbol_12(buffer, (offset + 0u)), tint_symbol_12(buffer, (offset + 16u)), tint_symbol_12(buffer, (offset + 32u)));
}
-fn tint_symbol_18([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat4x2<f32> {
- return mat4x2<f32>(tint_symbol_5(buffer, (offset + 0u)), tint_symbol_5(buffer, (offset + 8u)), tint_symbol_5(buffer, (offset + 16u)), tint_symbol_5(buffer, (offset + 24u)));
+fn tint_symbol_19([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat4x2<f32> {
+ return mat4x2<f32>(tint_symbol_6(buffer, (offset + 0u)), tint_symbol_6(buffer, (offset + 8u)), tint_symbol_6(buffer, (offset + 16u)), tint_symbol_6(buffer, (offset + 24u)));
}
-fn tint_symbol_19([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat4x3<f32> {
- return mat4x3<f32>(tint_symbol_8(buffer, (offset + 0u)), tint_symbol_8(buffer, (offset + 16u)), tint_symbol_8(buffer, (offset + 32u)), tint_symbol_8(buffer, (offset + 48u)));
+fn tint_symbol_20([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat4x3<f32> {
+ return mat4x3<f32>(tint_symbol_9(buffer, (offset + 0u)), tint_symbol_9(buffer, (offset + 16u)), tint_symbol_9(buffer, (offset + 32u)), tint_symbol_9(buffer, (offset + 48u)));
}
-fn tint_symbol_20([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat4x4<f32> {
- return mat4x4<f32>(tint_symbol_11(buffer, (offset + 0u)), tint_symbol_11(buffer, (offset + 16u)), tint_symbol_11(buffer, (offset + 32u)), tint_symbol_11(buffer, (offset + 48u)));
+fn tint_symbol_21([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat4x4<f32> {
+ return mat4x4<f32>(tint_symbol_12(buffer, (offset + 0u)), tint_symbol_12(buffer, (offset + 16u)), tint_symbol_12(buffer, (offset + 32u)), tint_symbol_12(buffer, (offset + 48u)));
}
-fn tint_symbol_21([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> array<vec3<f32>, 2> {
- return array<vec3<f32>, 2>(tint_symbol_8(buffer, (offset + 0u)), tint_symbol_8(buffer, (offset + 16u)));
+fn tint_symbol_22([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> array<vec3<f32>, 2> {
+ var arr : array<vec3<f32>, 2>;
+ for(var i_1 = 0u; (i_1 < 2u); i_1 = (i_1 + 1u)) {
+ arr[i_1] = tint_symbol_9(buffer, (offset + (i_1 * 16u)));
+ }
+ return arr;
}
-fn tint_symbol_22([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> SB {
- return SB(tint_symbol(buffer, (offset + 0u)), tint_symbol_1(buffer, (offset + 4u)), tint_symbol_2(buffer, (offset + 8u)), tint_symbol_3(buffer, (offset + 16u)), tint_symbol_4(buffer, (offset + 24u)), tint_symbol_5(buffer, (offset + 32u)), tint_symbol_6(buffer, (offset + 48u)), tint_symbol_7(buffer, (offset + 64u)), tint_symbol_8(buffer, (offset + 80u)), tint_symbol_9(buffer, (offset + 96u)), tint_symbol_10(buffer, (offset + 112u)), tint_symbol_11(buffer, (offset + 128u)), tint_symbol_12(buffer, (offset + 144u)), tint_symbol_13(buffer, (offset + 160u)), tint_symbol_14(buffer, (offset + 192u)), tint_symbol_15(buffer, (offset + 224u)), tint_symbol_16(buffer, (offset + 256u)), tint_symbol_17(buffer, (offset + 304u)), tint_symbol_18(buffer, (offset + 352u)), tint_symbol_19(buffer, (offset + 384u)), tint_symbol_20(buffer, (offset + 448u)), tint_symbol_21(buffer, (offset + 512u)));
+fn tint_symbol([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> SB {
+ return SB(tint_symbol_1(buffer, (offset + 0u)), tint_symbol_2(buffer, (offset + 4u)), tint_symbol_3(buffer, (offset + 8u)), tint_symbol_4(buffer, (offset + 16u)), tint_symbol_5(buffer, (offset + 24u)), tint_symbol_6(buffer, (offset + 32u)), tint_symbol_7(buffer, (offset + 48u)), tint_symbol_8(buffer, (offset + 64u)), tint_symbol_9(buffer, (offset + 80u)), tint_symbol_10(buffer, (offset + 96u)), tint_symbol_11(buffer, (offset + 112u)), tint_symbol_12(buffer, (offset + 128u)), tint_symbol_13(buffer, (offset + 144u)), tint_symbol_14(buffer, (offset + 160u)), tint_symbol_15(buffer, (offset + 192u)), tint_symbol_16(buffer, (offset + 224u)), tint_symbol_17(buffer, (offset + 256u)), tint_symbol_18(buffer, (offset + 304u)), tint_symbol_19(buffer, (offset + 352u)), tint_symbol_20(buffer, (offset + 384u)), tint_symbol_21(buffer, (offset + 448u)), tint_symbol_22(buffer, (offset + 512u)));
}
[[stage(compute), workgroup_size(1)]]
fn main() {
- var x : SB = tint_symbol_22(sb, 0u);
+ var x : SB = tint_symbol(sb, 0u);
}
)";
@@ -845,128 +859,130 @@
[[group(0), binding(0)]] var<storage, read_write> sb : SB;
[[internal(intrinsic_store_storage_i32), internal(disable_validation__function_has_no_body)]]
-fn tint_symbol([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : i32)
+fn tint_symbol_1([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : i32)
[[internal(intrinsic_store_storage_u32), internal(disable_validation__function_has_no_body)]]
-fn tint_symbol_1([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : u32)
+fn tint_symbol_2([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : u32)
[[internal(intrinsic_store_storage_f32), internal(disable_validation__function_has_no_body)]]
-fn tint_symbol_2([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : f32)
+fn tint_symbol_3([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : f32)
[[internal(intrinsic_store_storage_vec2_i32), internal(disable_validation__function_has_no_body)]]
-fn tint_symbol_3([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec2<i32>)
+fn tint_symbol_4([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec2<i32>)
[[internal(intrinsic_store_storage_vec2_u32), internal(disable_validation__function_has_no_body)]]
-fn tint_symbol_4([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec2<u32>)
+fn tint_symbol_5([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec2<u32>)
[[internal(intrinsic_store_storage_vec2_f32), internal(disable_validation__function_has_no_body)]]
-fn tint_symbol_5([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec2<f32>)
+fn tint_symbol_6([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec2<f32>)
[[internal(intrinsic_store_storage_vec3_i32), internal(disable_validation__function_has_no_body)]]
-fn tint_symbol_6([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec3<i32>)
+fn tint_symbol_7([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec3<i32>)
[[internal(intrinsic_store_storage_vec3_u32), internal(disable_validation__function_has_no_body)]]
-fn tint_symbol_7([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec3<u32>)
+fn tint_symbol_8([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec3<u32>)
[[internal(intrinsic_store_storage_vec3_f32), internal(disable_validation__function_has_no_body)]]
-fn tint_symbol_8([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec3<f32>)
+fn tint_symbol_9([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec3<f32>)
[[internal(intrinsic_store_storage_vec4_i32), internal(disable_validation__function_has_no_body)]]
-fn tint_symbol_9([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec4<i32>)
+fn tint_symbol_10([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec4<i32>)
[[internal(intrinsic_store_storage_vec4_u32), internal(disable_validation__function_has_no_body)]]
-fn tint_symbol_10([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec4<u32>)
+fn tint_symbol_11([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec4<u32>)
[[internal(intrinsic_store_storage_vec4_f32), internal(disable_validation__function_has_no_body)]]
-fn tint_symbol_11([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec4<f32>)
+fn tint_symbol_12([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec4<f32>)
-fn tint_symbol_12([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat2x2<f32>) {
- tint_symbol_5(buffer, (offset + 0u), value[0u]);
- tint_symbol_5(buffer, (offset + 8u), value[1u]);
+fn tint_symbol_13([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat2x2<f32>) {
+ tint_symbol_6(buffer, (offset + 0u), value[0u]);
+ tint_symbol_6(buffer, (offset + 8u), value[1u]);
}
-fn tint_symbol_13([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat2x3<f32>) {
- tint_symbol_8(buffer, (offset + 0u), value[0u]);
- tint_symbol_8(buffer, (offset + 16u), value[1u]);
+fn tint_symbol_14([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat2x3<f32>) {
+ tint_symbol_9(buffer, (offset + 0u), value[0u]);
+ tint_symbol_9(buffer, (offset + 16u), value[1u]);
}
-fn tint_symbol_14([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat2x4<f32>) {
- tint_symbol_11(buffer, (offset + 0u), value[0u]);
- tint_symbol_11(buffer, (offset + 16u), value[1u]);
+fn tint_symbol_15([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat2x4<f32>) {
+ tint_symbol_12(buffer, (offset + 0u), value[0u]);
+ tint_symbol_12(buffer, (offset + 16u), value[1u]);
}
-fn tint_symbol_15([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat3x2<f32>) {
- tint_symbol_5(buffer, (offset + 0u), value[0u]);
- tint_symbol_5(buffer, (offset + 8u), value[1u]);
- tint_symbol_5(buffer, (offset + 16u), value[2u]);
+fn tint_symbol_16([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat3x2<f32>) {
+ tint_symbol_6(buffer, (offset + 0u), value[0u]);
+ tint_symbol_6(buffer, (offset + 8u), value[1u]);
+ tint_symbol_6(buffer, (offset + 16u), value[2u]);
}
-fn tint_symbol_16([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat3x3<f32>) {
- tint_symbol_8(buffer, (offset + 0u), value[0u]);
- tint_symbol_8(buffer, (offset + 16u), value[1u]);
- tint_symbol_8(buffer, (offset + 32u), value[2u]);
+fn tint_symbol_17([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat3x3<f32>) {
+ tint_symbol_9(buffer, (offset + 0u), value[0u]);
+ tint_symbol_9(buffer, (offset + 16u), value[1u]);
+ tint_symbol_9(buffer, (offset + 32u), value[2u]);
}
-fn tint_symbol_17([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat3x4<f32>) {
- tint_symbol_11(buffer, (offset + 0u), value[0u]);
- tint_symbol_11(buffer, (offset + 16u), value[1u]);
- tint_symbol_11(buffer, (offset + 32u), value[2u]);
+fn tint_symbol_18([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat3x4<f32>) {
+ tint_symbol_12(buffer, (offset + 0u), value[0u]);
+ tint_symbol_12(buffer, (offset + 16u), value[1u]);
+ tint_symbol_12(buffer, (offset + 32u), value[2u]);
}
-fn tint_symbol_18([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat4x2<f32>) {
- tint_symbol_5(buffer, (offset + 0u), value[0u]);
- tint_symbol_5(buffer, (offset + 8u), value[1u]);
- tint_symbol_5(buffer, (offset + 16u), value[2u]);
- tint_symbol_5(buffer, (offset + 24u), value[3u]);
+fn tint_symbol_19([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat4x2<f32>) {
+ tint_symbol_6(buffer, (offset + 0u), value[0u]);
+ tint_symbol_6(buffer, (offset + 8u), value[1u]);
+ tint_symbol_6(buffer, (offset + 16u), value[2u]);
+ tint_symbol_6(buffer, (offset + 24u), value[3u]);
}
-fn tint_symbol_19([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat4x3<f32>) {
- tint_symbol_8(buffer, (offset + 0u), value[0u]);
- tint_symbol_8(buffer, (offset + 16u), value[1u]);
- tint_symbol_8(buffer, (offset + 32u), value[2u]);
- tint_symbol_8(buffer, (offset + 48u), value[3u]);
+fn tint_symbol_20([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat4x3<f32>) {
+ tint_symbol_9(buffer, (offset + 0u), value[0u]);
+ tint_symbol_9(buffer, (offset + 16u), value[1u]);
+ tint_symbol_9(buffer, (offset + 32u), value[2u]);
+ tint_symbol_9(buffer, (offset + 48u), value[3u]);
}
-fn tint_symbol_20([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat4x4<f32>) {
- tint_symbol_11(buffer, (offset + 0u), value[0u]);
- tint_symbol_11(buffer, (offset + 16u), value[1u]);
- tint_symbol_11(buffer, (offset + 32u), value[2u]);
- tint_symbol_11(buffer, (offset + 48u), value[3u]);
+fn tint_symbol_21([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat4x4<f32>) {
+ tint_symbol_12(buffer, (offset + 0u), value[0u]);
+ tint_symbol_12(buffer, (offset + 16u), value[1u]);
+ tint_symbol_12(buffer, (offset + 32u), value[2u]);
+ tint_symbol_12(buffer, (offset + 48u), value[3u]);
}
-fn tint_symbol_21([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : array<vec3<f32>, 2>) {
- tint_symbol_8(buffer, (offset + 0u), value[0u]);
- tint_symbol_8(buffer, (offset + 16u), value[1u]);
+fn tint_symbol_22([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : array<vec3<f32>, 2>) {
+ var array = value;
+ for(var i_1 = 0u; (i_1 < 2u); i_1 = (i_1 + 1u)) {
+ tint_symbol_9(buffer, (offset + (i_1 * 16u)), array[i_1]);
+ }
}
-fn tint_symbol_22([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : SB) {
- tint_symbol(buffer, (offset + 0u), value.a);
- tint_symbol_1(buffer, (offset + 4u), value.b);
- tint_symbol_2(buffer, (offset + 8u), value.c);
- tint_symbol_3(buffer, (offset + 16u), value.d);
- tint_symbol_4(buffer, (offset + 24u), value.e);
- tint_symbol_5(buffer, (offset + 32u), value.f);
- tint_symbol_6(buffer, (offset + 48u), value.g);
- tint_symbol_7(buffer, (offset + 64u), value.h);
- tint_symbol_8(buffer, (offset + 80u), value.i);
- tint_symbol_9(buffer, (offset + 96u), value.j);
- tint_symbol_10(buffer, (offset + 112u), value.k);
- tint_symbol_11(buffer, (offset + 128u), value.l);
- tint_symbol_12(buffer, (offset + 144u), value.m);
- tint_symbol_13(buffer, (offset + 160u), value.n);
- tint_symbol_14(buffer, (offset + 192u), value.o);
- tint_symbol_15(buffer, (offset + 224u), value.p);
- tint_symbol_16(buffer, (offset + 256u), value.q);
- tint_symbol_17(buffer, (offset + 304u), value.r);
- tint_symbol_18(buffer, (offset + 352u), value.s);
- tint_symbol_19(buffer, (offset + 384u), value.t);
- tint_symbol_20(buffer, (offset + 448u), value.u);
- tint_symbol_21(buffer, (offset + 512u), value.v);
+fn tint_symbol([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : SB) {
+ tint_symbol_1(buffer, (offset + 0u), value.a);
+ tint_symbol_2(buffer, (offset + 4u), value.b);
+ tint_symbol_3(buffer, (offset + 8u), value.c);
+ tint_symbol_4(buffer, (offset + 16u), value.d);
+ tint_symbol_5(buffer, (offset + 24u), value.e);
+ tint_symbol_6(buffer, (offset + 32u), value.f);
+ tint_symbol_7(buffer, (offset + 48u), value.g);
+ tint_symbol_8(buffer, (offset + 64u), value.h);
+ tint_symbol_9(buffer, (offset + 80u), value.i);
+ tint_symbol_10(buffer, (offset + 96u), value.j);
+ tint_symbol_11(buffer, (offset + 112u), value.k);
+ tint_symbol_12(buffer, (offset + 128u), value.l);
+ tint_symbol_13(buffer, (offset + 144u), value.m);
+ tint_symbol_14(buffer, (offset + 160u), value.n);
+ tint_symbol_15(buffer, (offset + 192u), value.o);
+ tint_symbol_16(buffer, (offset + 224u), value.p);
+ tint_symbol_17(buffer, (offset + 256u), value.q);
+ tint_symbol_18(buffer, (offset + 304u), value.r);
+ tint_symbol_19(buffer, (offset + 352u), value.s);
+ tint_symbol_20(buffer, (offset + 384u), value.t);
+ tint_symbol_21(buffer, (offset + 448u), value.u);
+ tint_symbol_22(buffer, (offset + 512u), value.v);
}
[[stage(compute), workgroup_size(1)]]
fn main() {
- tint_symbol_22(sb, 0u, SB());
+ tint_symbol(sb, 0u, SB());
}
)";
diff --git a/src/writer/hlsl/generator_impl_member_accessor_test.cc b/src/writer/hlsl/generator_impl_member_accessor_test.cc
index 10d7cb9..40b73d4 100644
--- a/src/writer/hlsl/generator_impl_member_accessor_test.cc
+++ b/src/writer/hlsl/generator_impl_member_accessor_test.cc
@@ -343,13 +343,13 @@
auto* expected =
R"(RWByteAddressBuffer data : register(u0, space1);
-void tint_symbol_1(RWByteAddressBuffer buffer, uint offset, float2x3 value) {
+void tint_symbol(RWByteAddressBuffer buffer, uint offset, float2x3 value) {
buffer.Store3((offset + 0u), asuint(value[0u]));
buffer.Store3((offset + 16u), asuint(value[1u]));
}
void main() {
- tint_symbol_1(data, 16u, float2x3(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f));
+ tint_symbol(data, 16u, float2x3(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f));
return;
}
)";
diff --git a/test/array/assign_to_function_var.wgsl.expected.hlsl b/test/array/assign_to_function_var.wgsl.expected.hlsl
index 2c7d6be..5b15fc1 100644
--- a/test/array/assign_to_function_var.wgsl.expected.hlsl
+++ b/test/array/assign_to_function_var.wgsl.expected.hlsl
@@ -28,27 +28,34 @@
return tint_symbol_6;
}
-typedef tint_padded_array_element tint_symbol_2_ret[4];
-tint_symbol_2_ret tint_symbol_2(uint4 buffer[4], uint offset) {
- const uint scalar_offset = ((offset + 0u)) / 4;
- const uint scalar_offset_1 = ((offset + 16u)) / 4;
- const uint scalar_offset_2 = ((offset + 32u)) / 4;
- const uint scalar_offset_3 = ((offset + 48u)) / 4;
- const tint_padded_array_element tint_symbol_7[4] = {{asint(buffer[scalar_offset / 4][scalar_offset % 4])}, {asint(buffer[scalar_offset_1 / 4][scalar_offset_1 % 4])}, {asint(buffer[scalar_offset_2 / 4][scalar_offset_2 % 4])}, {asint(buffer[scalar_offset_3 / 4][scalar_offset_3 % 4])}};
- return tint_symbol_7;
+typedef tint_padded_array_element tint_symbol_1_ret[4];
+tint_symbol_1_ret tint_symbol_1(uint4 buffer[4], uint offset) {
+ tint_padded_array_element arr_1[4] = (tint_padded_array_element[4])0;
+ {
+ for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+ const uint scalar_offset = ((offset + (i * 16u))) / 4;
+ arr_1[i].el = asint(buffer[scalar_offset / 4][scalar_offset % 4]);
+ }
+ }
+ return arr_1;
}
-typedef tint_padded_array_element tint_symbol_4_ret[4];
-tint_symbol_4_ret tint_symbol_4(RWByteAddressBuffer buffer, uint offset) {
- const tint_padded_array_element tint_symbol_8[4] = {{asint(buffer.Load((offset + 0u)))}, {asint(buffer.Load((offset + 16u)))}, {asint(buffer.Load((offset + 32u)))}, {asint(buffer.Load((offset + 48u)))}};
- return tint_symbol_8;
+typedef tint_padded_array_element tint_symbol_3_ret[4];
+tint_symbol_3_ret tint_symbol_3(RWByteAddressBuffer buffer, uint offset) {
+ tint_padded_array_element arr_2[4] = (tint_padded_array_element[4])0;
+ {
+ for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
+ arr_2[i_1].el = asint(buffer.Load((offset + (i_1 * 16u))));
+ }
+ }
+ return arr_2;
}
void foo(tint_padded_array_element src_param[4]) {
tint_padded_array_element src_function[4] = (tint_padded_array_element[4])0;
tint_padded_array_element tint_symbol[4] = (tint_padded_array_element[4])0;
- const tint_padded_array_element tint_symbol_9[4] = {{1}, {2}, {3}, {3}};
- tint_symbol = tint_symbol_9;
+ const tint_padded_array_element tint_symbol_7[4] = {{1}, {2}, {3}, {3}};
+ tint_symbol = tint_symbol_7;
tint_symbol = src_param;
tint_symbol = ret_arr();
const tint_padded_array_element src_let[4] = (tint_padded_array_element[4])0;
@@ -57,8 +64,8 @@
tint_symbol = src_private;
tint_symbol = src_workgroup;
tint_symbol = ret_struct_arr().arr;
- tint_symbol = tint_symbol_2(src_uniform, 0u);
- tint_symbol = tint_symbol_4(src_storage, 0u);
+ tint_symbol = tint_symbol_1(src_uniform, 0u);
+ tint_symbol = tint_symbol_3(src_storage, 0u);
int dst_nested[4][3][2] = (int[4][3][2])0;
int src_nested[4][3][2] = (int[4][3][2])0;
dst_nested = src_nested;
diff --git a/test/array/assign_to_private_var.wgsl.expected.hlsl b/test/array/assign_to_private_var.wgsl.expected.hlsl
index 68ea5c3..8cd8ee1 100644
--- a/test/array/assign_to_private_var.wgsl.expected.hlsl
+++ b/test/array/assign_to_private_var.wgsl.expected.hlsl
@@ -30,26 +30,33 @@
return tint_symbol_6;
}
-typedef tint_padded_array_element tint_symbol_2_ret[4];
-tint_symbol_2_ret tint_symbol_2(uint4 buffer[4], uint offset) {
- const uint scalar_offset = ((offset + 0u)) / 4;
- const uint scalar_offset_1 = ((offset + 16u)) / 4;
- const uint scalar_offset_2 = ((offset + 32u)) / 4;
- const uint scalar_offset_3 = ((offset + 48u)) / 4;
- const tint_padded_array_element tint_symbol_7[4] = {{asint(buffer[scalar_offset / 4][scalar_offset % 4])}, {asint(buffer[scalar_offset_1 / 4][scalar_offset_1 % 4])}, {asint(buffer[scalar_offset_2 / 4][scalar_offset_2 % 4])}, {asint(buffer[scalar_offset_3 / 4][scalar_offset_3 % 4])}};
- return tint_symbol_7;
+typedef tint_padded_array_element tint_symbol_1_ret[4];
+tint_symbol_1_ret tint_symbol_1(uint4 buffer[4], uint offset) {
+ tint_padded_array_element arr_1[4] = (tint_padded_array_element[4])0;
+ {
+ for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+ const uint scalar_offset = ((offset + (i * 16u))) / 4;
+ arr_1[i].el = asint(buffer[scalar_offset / 4][scalar_offset % 4]);
+ }
+ }
+ return arr_1;
}
-typedef tint_padded_array_element tint_symbol_4_ret[4];
-tint_symbol_4_ret tint_symbol_4(RWByteAddressBuffer buffer, uint offset) {
- const tint_padded_array_element tint_symbol_8[4] = {{asint(buffer.Load((offset + 0u)))}, {asint(buffer.Load((offset + 16u)))}, {asint(buffer.Load((offset + 32u)))}, {asint(buffer.Load((offset + 48u)))}};
- return tint_symbol_8;
+typedef tint_padded_array_element tint_symbol_3_ret[4];
+tint_symbol_3_ret tint_symbol_3(RWByteAddressBuffer buffer, uint offset) {
+ tint_padded_array_element arr_2[4] = (tint_padded_array_element[4])0;
+ {
+ for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
+ arr_2[i_1].el = asint(buffer.Load((offset + (i_1 * 16u))));
+ }
+ }
+ return arr_2;
}
void foo(tint_padded_array_element src_param[4]) {
tint_padded_array_element src_function[4] = (tint_padded_array_element[4])0;
- const tint_padded_array_element tint_symbol_9[4] = {{1}, {2}, {3}, {3}};
- tint_symbol = tint_symbol_9;
+ const tint_padded_array_element tint_symbol_7[4] = {{1}, {2}, {3}, {3}};
+ tint_symbol = tint_symbol_7;
tint_symbol = src_param;
tint_symbol = ret_arr();
const tint_padded_array_element src_let[4] = (tint_padded_array_element[4])0;
@@ -58,8 +65,8 @@
tint_symbol = src_private;
tint_symbol = src_workgroup;
tint_symbol = ret_struct_arr().arr;
- tint_symbol = tint_symbol_2(src_uniform, 0u);
- tint_symbol = tint_symbol_4(src_storage, 0u);
+ tint_symbol = tint_symbol_1(src_uniform, 0u);
+ tint_symbol = tint_symbol_3(src_storage, 0u);
int src_nested[4][3][2] = (int[4][3][2])0;
dst_nested = src_nested;
}
diff --git a/test/array/assign_to_storage_var.wgsl.expected.hlsl b/test/array/assign_to_storage_var.wgsl.expected.hlsl
index ee58066..8825ef3 100644
--- a/test/array/assign_to_storage_var.wgsl.expected.hlsl
+++ b/test/array/assign_to_storage_var.wgsl.expected.hlsl
@@ -30,61 +30,79 @@
return tint_symbol_12;
}
-void tint_symbol_2(RWByteAddressBuffer buffer, uint offset, tint_padded_array_element value[4]) {
- buffer.Store((offset + 0u), asuint(value[0u].el));
- buffer.Store((offset + 16u), asuint(value[1u].el));
- buffer.Store((offset + 32u), asuint(value[2u].el));
- buffer.Store((offset + 48u), asuint(value[3u].el));
+void tint_symbol_1(RWByteAddressBuffer buffer, uint offset, tint_padded_array_element value[4]) {
+ tint_padded_array_element array[4] = value;
+ {
+ for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+ buffer.Store((offset + (i * 16u)), asuint(array[i].el));
+ }
+ }
}
-typedef tint_padded_array_element tint_symbol_4_ret[4];
-tint_symbol_4_ret tint_symbol_4(uint4 buffer[4], uint offset) {
- const uint scalar_offset = ((offset + 0u)) / 4;
- const uint scalar_offset_1 = ((offset + 16u)) / 4;
- const uint scalar_offset_2 = ((offset + 32u)) / 4;
- const uint scalar_offset_3 = ((offset + 48u)) / 4;
- const tint_padded_array_element tint_symbol_13[4] = {{asint(buffer[scalar_offset / 4][scalar_offset % 4])}, {asint(buffer[scalar_offset_1 / 4][scalar_offset_1 % 4])}, {asint(buffer[scalar_offset_2 / 4][scalar_offset_2 % 4])}, {asint(buffer[scalar_offset_3 / 4][scalar_offset_3 % 4])}};
- return tint_symbol_13;
+typedef tint_padded_array_element tint_symbol_3_ret[4];
+tint_symbol_3_ret tint_symbol_3(uint4 buffer[4], uint offset) {
+ tint_padded_array_element arr_1[4] = (tint_padded_array_element[4])0;
+ {
+ for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
+ const uint scalar_offset = ((offset + (i_1 * 16u))) / 4;
+ arr_1[i_1].el = asint(buffer[scalar_offset / 4][scalar_offset % 4]);
+ }
+ }
+ return arr_1;
}
-typedef tint_padded_array_element tint_symbol_6_ret[4];
-tint_symbol_6_ret tint_symbol_6(RWByteAddressBuffer buffer, uint offset) {
- const tint_padded_array_element tint_symbol_14[4] = {{asint(buffer.Load((offset + 0u)))}, {asint(buffer.Load((offset + 16u)))}, {asint(buffer.Load((offset + 32u)))}, {asint(buffer.Load((offset + 48u)))}};
- return tint_symbol_14;
+typedef tint_padded_array_element tint_symbol_5_ret[4];
+tint_symbol_5_ret tint_symbol_5(RWByteAddressBuffer buffer, uint offset) {
+ tint_padded_array_element arr_2[4] = (tint_padded_array_element[4])0;
+ {
+ for(uint i_2 = 0u; (i_2 < 4u); i_2 = (i_2 + 1u)) {
+ arr_2[i_2].el = asint(buffer.Load((offset + (i_2 * 16u))));
+ }
+ }
+ return arr_2;
}
-void tint_symbol_8(RWByteAddressBuffer buffer, uint offset, int value[2]) {
- buffer.Store((offset + 0u), asuint(value[0u]));
- buffer.Store((offset + 4u), asuint(value[1u]));
+void tint_symbol_9(RWByteAddressBuffer buffer, uint offset, int value[2]) {
+ int array_3[2] = value;
+ {
+ for(uint i_3 = 0u; (i_3 < 2u); i_3 = (i_3 + 1u)) {
+ buffer.Store((offset + (i_3 * 4u)), asuint(array_3[i_3]));
+ }
+ }
}
-void tint_symbol_9(RWByteAddressBuffer buffer, uint offset, int value[3][2]) {
- tint_symbol_8(buffer, (offset + 0u), value[0u]);
- tint_symbol_8(buffer, (offset + 8u), value[1u]);
- tint_symbol_8(buffer, (offset + 16u), value[2u]);
+void tint_symbol_8(RWByteAddressBuffer buffer, uint offset, int value[3][2]) {
+ int array_2[3][2] = value;
+ {
+ for(uint i_4 = 0u; (i_4 < 3u); i_4 = (i_4 + 1u)) {
+ tint_symbol_9(buffer, (offset + (i_4 * 8u)), array_2[i_4]);
+ }
+ }
}
-void tint_symbol_10(RWByteAddressBuffer buffer, uint offset, int value[4][3][2]) {
- tint_symbol_9(buffer, (offset + 0u), value[0u]);
- tint_symbol_9(buffer, (offset + 24u), value[1u]);
- tint_symbol_9(buffer, (offset + 48u), value[2u]);
- tint_symbol_9(buffer, (offset + 72u), value[3u]);
+void tint_symbol_7(RWByteAddressBuffer buffer, uint offset, int value[4][3][2]) {
+ int array_1[4][3][2] = value;
+ {
+ for(uint i_5 = 0u; (i_5 < 4u); i_5 = (i_5 + 1u)) {
+ tint_symbol_8(buffer, (offset + (i_5 * 24u)), array_1[i_5]);
+ }
+ }
}
void foo(tint_padded_array_element src_param[4]) {
tint_padded_array_element src_function[4] = (tint_padded_array_element[4])0;
- const tint_padded_array_element tint_symbol_15[4] = {{1}, {2}, {3}, {3}};
- tint_symbol_2(tint_symbol, 0u, tint_symbol_15);
- tint_symbol_2(tint_symbol, 0u, src_param);
- tint_symbol_2(tint_symbol, 0u, ret_arr());
+ const tint_padded_array_element tint_symbol_13[4] = {{1}, {2}, {3}, {3}};
+ tint_symbol_1(tint_symbol, 0u, tint_symbol_13);
+ tint_symbol_1(tint_symbol, 0u, src_param);
+ tint_symbol_1(tint_symbol, 0u, ret_arr());
const tint_padded_array_element src_let[4] = (tint_padded_array_element[4])0;
- tint_symbol_2(tint_symbol, 0u, src_let);
- tint_symbol_2(tint_symbol, 0u, src_function);
- tint_symbol_2(tint_symbol, 0u, src_private);
- tint_symbol_2(tint_symbol, 0u, src_workgroup);
- tint_symbol_2(tint_symbol, 0u, ret_struct_arr().arr);
- tint_symbol_2(tint_symbol, 0u, tint_symbol_4(src_uniform, 0u));
- tint_symbol_2(tint_symbol, 0u, tint_symbol_6(src_storage, 0u));
+ tint_symbol_1(tint_symbol, 0u, src_let);
+ tint_symbol_1(tint_symbol, 0u, src_function);
+ tint_symbol_1(tint_symbol, 0u, src_private);
+ tint_symbol_1(tint_symbol, 0u, src_workgroup);
+ tint_symbol_1(tint_symbol, 0u, ret_struct_arr().arr);
+ tint_symbol_1(tint_symbol, 0u, tint_symbol_3(src_uniform, 0u));
+ tint_symbol_1(tint_symbol, 0u, tint_symbol_5(src_storage, 0u));
int src_nested[4][3][2] = (int[4][3][2])0;
- tint_symbol_10(dst_nested, 0u, src_nested);
+ tint_symbol_7(dst_nested, 0u, src_nested);
}
diff --git a/test/array/assign_to_workgroup_var.wgsl.expected.hlsl b/test/array/assign_to_workgroup_var.wgsl.expected.hlsl
index 48c5fae..43a8a1c 100644
--- a/test/array/assign_to_workgroup_var.wgsl.expected.hlsl
+++ b/test/array/assign_to_workgroup_var.wgsl.expected.hlsl
@@ -30,26 +30,33 @@
return tint_symbol_6;
}
-typedef tint_padded_array_element tint_symbol_2_ret[4];
-tint_symbol_2_ret tint_symbol_2(uint4 buffer[4], uint offset) {
- const uint scalar_offset = ((offset + 0u)) / 4;
- const uint scalar_offset_1 = ((offset + 16u)) / 4;
- const uint scalar_offset_2 = ((offset + 32u)) / 4;
- const uint scalar_offset_3 = ((offset + 48u)) / 4;
- const tint_padded_array_element tint_symbol_7[4] = {{asint(buffer[scalar_offset / 4][scalar_offset % 4])}, {asint(buffer[scalar_offset_1 / 4][scalar_offset_1 % 4])}, {asint(buffer[scalar_offset_2 / 4][scalar_offset_2 % 4])}, {asint(buffer[scalar_offset_3 / 4][scalar_offset_3 % 4])}};
- return tint_symbol_7;
+typedef tint_padded_array_element tint_symbol_1_ret[4];
+tint_symbol_1_ret tint_symbol_1(uint4 buffer[4], uint offset) {
+ tint_padded_array_element arr_1[4] = (tint_padded_array_element[4])0;
+ {
+ for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+ const uint scalar_offset = ((offset + (i * 16u))) / 4;
+ arr_1[i].el = asint(buffer[scalar_offset / 4][scalar_offset % 4]);
+ }
+ }
+ return arr_1;
}
-typedef tint_padded_array_element tint_symbol_4_ret[4];
-tint_symbol_4_ret tint_symbol_4(RWByteAddressBuffer buffer, uint offset) {
- const tint_padded_array_element tint_symbol_8[4] = {{asint(buffer.Load((offset + 0u)))}, {asint(buffer.Load((offset + 16u)))}, {asint(buffer.Load((offset + 32u)))}, {asint(buffer.Load((offset + 48u)))}};
- return tint_symbol_8;
+typedef tint_padded_array_element tint_symbol_3_ret[4];
+tint_symbol_3_ret tint_symbol_3(RWByteAddressBuffer buffer, uint offset) {
+ tint_padded_array_element arr_2[4] = (tint_padded_array_element[4])0;
+ {
+ for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
+ arr_2[i_1].el = asint(buffer.Load((offset + (i_1 * 16u))));
+ }
+ }
+ return arr_2;
}
void foo(tint_padded_array_element src_param[4]) {
tint_padded_array_element src_function[4] = (tint_padded_array_element[4])0;
- const tint_padded_array_element tint_symbol_9[4] = {{1}, {2}, {3}, {3}};
- tint_symbol = tint_symbol_9;
+ const tint_padded_array_element tint_symbol_7[4] = {{1}, {2}, {3}, {3}};
+ tint_symbol = tint_symbol_7;
tint_symbol = src_param;
tint_symbol = ret_arr();
const tint_padded_array_element src_let[4] = (tint_padded_array_element[4])0;
@@ -58,8 +65,8 @@
tint_symbol = src_private;
tint_symbol = src_workgroup;
tint_symbol = ret_struct_arr().arr;
- tint_symbol = tint_symbol_2(src_uniform, 0u);
- tint_symbol = tint_symbol_4(src_storage, 0u);
+ tint_symbol = tint_symbol_1(src_uniform, 0u);
+ tint_symbol = tint_symbol_3(src_storage, 0u);
int src_nested[4][3][2] = (int[4][3][2])0;
dst_nested = src_nested;
}
diff --git a/test/buffer/storage/dynamic_index/read.wgsl.expected.hlsl b/test/buffer/storage/dynamic_index/read.wgsl.expected.hlsl
index 9b9390a..24312d1 100644
--- a/test/buffer/storage/dynamic_index/read.wgsl.expected.hlsl
+++ b/test/buffer/storage/dynamic_index/read.wgsl.expected.hlsl
@@ -8,14 +8,19 @@
return float2x3(asfloat(buffer.Load3((offset + 0u))), asfloat(buffer.Load3((offset + 16u))));
}
-float3x2 tint_symbol_10(ByteAddressBuffer buffer, uint offset) {
+float3x2 tint_symbol_9(ByteAddressBuffer buffer, uint offset) {
return float3x2(asfloat(buffer.Load2((offset + 0u))), asfloat(buffer.Load2((offset + 8u))), asfloat(buffer.Load2((offset + 16u))));
}
-typedef int4 tint_symbol_12_ret[4];
-tint_symbol_12_ret tint_symbol_12(ByteAddressBuffer buffer, uint offset) {
- const int4 tint_symbol_13[4] = {asint(buffer.Load4((offset + 0u))), asint(buffer.Load4((offset + 16u))), asint(buffer.Load4((offset + 32u))), asint(buffer.Load4((offset + 48u)))};
- return tint_symbol_13;
+typedef int4 tint_symbol_11_ret[4];
+tint_symbol_11_ret tint_symbol_11(ByteAddressBuffer buffer, uint offset) {
+ int4 arr_1[4] = (int4[4])0;
+ {
+ for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
+ arr_1[i_1] = asint(buffer.Load4((offset + (i_1 * 16u))));
+ }
+ }
+ return arr_1;
}
[numthreads(1, 1, 1)]
@@ -28,7 +33,7 @@
const float3 e = asfloat(s.Load3(((176u * idx) + 32u)));
const float f = asfloat(s.Load(((176u * idx) + 44u)));
const float2x3 g = tint_symbol_8(s, ((176u * idx) + 48u));
- const float3x2 h = tint_symbol_10(s, ((176u * idx) + 80u));
- const int4 i[4] = tint_symbol_12(s, ((176u * idx) + 112u));
+ const float3x2 h = tint_symbol_9(s, ((176u * idx) + 80u));
+ const int4 i[4] = tint_symbol_11(s, ((176u * idx) + 112u));
return;
}
diff --git a/test/buffer/storage/dynamic_index/write.wgsl.expected.hlsl b/test/buffer/storage/dynamic_index/write.wgsl.expected.hlsl
index 74aeaf2..c18b7c6 100644
--- a/test/buffer/storage/dynamic_index/write.wgsl.expected.hlsl
+++ b/test/buffer/storage/dynamic_index/write.wgsl.expected.hlsl
@@ -9,17 +9,19 @@
buffer.Store3((offset + 16u), asuint(value[1u]));
}
-void tint_symbol_10(RWByteAddressBuffer buffer, uint offset, float3x2 value) {
+void tint_symbol_9(RWByteAddressBuffer buffer, uint offset, float3x2 value) {
buffer.Store2((offset + 0u), asuint(value[0u]));
buffer.Store2((offset + 8u), asuint(value[1u]));
buffer.Store2((offset + 16u), asuint(value[2u]));
}
-void tint_symbol_12(RWByteAddressBuffer buffer, uint offset, int4 value[4]) {
- buffer.Store4((offset + 0u), asuint(value[0u]));
- buffer.Store4((offset + 16u), asuint(value[1u]));
- buffer.Store4((offset + 32u), asuint(value[2u]));
- buffer.Store4((offset + 48u), asuint(value[3u]));
+void tint_symbol_11(RWByteAddressBuffer buffer, uint offset, int4 value[4]) {
+ int4 array[4] = value;
+ {
+ for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
+ buffer.Store4((offset + (i_1 * 16u)), asuint(array[i_1]));
+ }
+ }
}
[numthreads(1, 1, 1)]
@@ -32,8 +34,8 @@
s.Store3(((176u * idx) + 32u), asuint(float3(0.0f, 0.0f, 0.0f)));
s.Store(((176u * idx) + 44u), asuint(0.0f));
tint_symbol_8(s, ((176u * idx) + 48u), float2x3(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f));
- tint_symbol_10(s, ((176u * idx) + 80u), float3x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f));
+ tint_symbol_9(s, ((176u * idx) + 80u), float3x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f));
const int4 tint_symbol_13[4] = (int4[4])0;
- tint_symbol_12(s, ((176u * idx) + 112u), tint_symbol_13);
+ tint_symbol_11(s, ((176u * idx) + 112u), tint_symbol_13);
return;
}
diff --git a/test/buffer/storage/static_index/read.wgsl.expected.hlsl b/test/buffer/storage/static_index/read.wgsl.expected.hlsl
index c14caff..42a6a71 100644
--- a/test/buffer/storage/static_index/read.wgsl.expected.hlsl
+++ b/test/buffer/storage/static_index/read.wgsl.expected.hlsl
@@ -11,7 +11,7 @@
return float2x3(asfloat(buffer.Load3((offset + 0u))), asfloat(buffer.Load3((offset + 16u))));
}
-float3x2 tint_symbol_8(ByteAddressBuffer buffer, uint offset) {
+float3x2 tint_symbol_7(ByteAddressBuffer buffer, uint offset) {
return float3x2(asfloat(buffer.Load2((offset + 0u))), asfloat(buffer.Load2((offset + 8u))), asfloat(buffer.Load2((offset + 16u))));
}
@@ -22,8 +22,13 @@
typedef tint_padded_array_element tint_symbol_10_ret[4];
tint_symbol_10_ret tint_symbol_10(ByteAddressBuffer buffer, uint offset) {
- const tint_padded_array_element tint_symbol_12[4] = {{tint_symbol_9(buffer, (offset + 0u))}, {tint_symbol_9(buffer, (offset + 16u))}, {tint_symbol_9(buffer, (offset + 32u))}, {tint_symbol_9(buffer, (offset + 48u))}};
- return tint_symbol_12;
+ tint_padded_array_element arr[4] = (tint_padded_array_element[4])0;
+ {
+ for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
+ arr[i_1].el = tint_symbol_9(buffer, (offset + (i_1 * 16u)));
+ }
+ }
+ return arr;
}
[numthreads(1, 1, 1)]
@@ -35,7 +40,7 @@
const float3 e = asfloat(s.Load3(32u));
const float f = asfloat(s.Load(44u));
const float2x3 g = tint_symbol_6(s, 48u);
- const float3x2 h = tint_symbol_8(s, 80u);
+ const float3x2 h = tint_symbol_7(s, 80u);
const Inner i = tint_symbol_9(s, 104u);
const tint_padded_array_element j[4] = tint_symbol_10(s, 108u);
return;
diff --git a/test/buffer/storage/static_index/write.wgsl.expected.hlsl b/test/buffer/storage/static_index/write.wgsl.expected.hlsl
index 7e569cd..8c9b7ea 100644
--- a/test/buffer/storage/static_index/write.wgsl.expected.hlsl
+++ b/test/buffer/storage/static_index/write.wgsl.expected.hlsl
@@ -12,7 +12,7 @@
buffer.Store3((offset + 16u), asuint(value[1u]));
}
-void tint_symbol_8(RWByteAddressBuffer buffer, uint offset, float3x2 value) {
+void tint_symbol_7(RWByteAddressBuffer buffer, uint offset, float3x2 value) {
buffer.Store2((offset + 0u), asuint(value[0u]));
buffer.Store2((offset + 8u), asuint(value[1u]));
buffer.Store2((offset + 16u), asuint(value[2u]));
@@ -23,10 +23,12 @@
}
void tint_symbol_10(RWByteAddressBuffer buffer, uint offset, tint_padded_array_element value[4]) {
- tint_symbol_9(buffer, (offset + 0u), value[0u].el);
- tint_symbol_9(buffer, (offset + 16u), value[1u].el);
- tint_symbol_9(buffer, (offset + 32u), value[2u].el);
- tint_symbol_9(buffer, (offset + 48u), value[3u].el);
+ tint_padded_array_element array[4] = value;
+ {
+ for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
+ tint_symbol_9(buffer, (offset + (i_1 * 16u)), array[i_1].el);
+ }
+ }
}
[numthreads(1, 1, 1)]
@@ -38,7 +40,7 @@
s.Store3(32u, asuint(float3(0.0f, 0.0f, 0.0f)));
s.Store(44u, asuint(0.0f));
tint_symbol_6(s, 48u, float2x3(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f));
- tint_symbol_8(s, 80u, float3x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f));
+ tint_symbol_7(s, 80u, float3x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f));
const Inner tint_symbol_11 = (Inner)0;
tint_symbol_9(s, 104u, tint_symbol_11);
const tint_padded_array_element tint_symbol_12[4] = (tint_padded_array_element[4])0;
diff --git a/test/buffer/uniform/static_index/read.wgsl.expected.hlsl b/test/buffer/uniform/static_index/read.wgsl.expected.hlsl
index fa6d133..f66c81b 100644
--- a/test/buffer/uniform/static_index/read.wgsl.expected.hlsl
+++ b/test/buffer/uniform/static_index/read.wgsl.expected.hlsl
@@ -15,7 +15,7 @@
return float2x3(asfloat(buffer[scalar_offset / 4].xyz), asfloat(buffer[scalar_offset_1 / 4].xyz));
}
-float3x2 tint_symbol_9(uint4 buffer[13], uint offset) {
+float3x2 tint_symbol_8(uint4 buffer[13], uint offset) {
const uint scalar_offset_2 = ((offset + 0u)) / 4;
uint4 ubo_load = buffer[scalar_offset_2 / 4];
const uint scalar_offset_3 = ((offset + 8u)) / 4;
@@ -33,8 +33,13 @@
typedef tint_padded_array_element tint_symbol_11_ret[4];
tint_symbol_11_ret tint_symbol_11(uint4 buffer[13], uint offset) {
- const tint_padded_array_element tint_symbol_13[4] = {{tint_symbol_10(buffer, (offset + 0u))}, {tint_symbol_10(buffer, (offset + 16u))}, {tint_symbol_10(buffer, (offset + 32u))}, {tint_symbol_10(buffer, (offset + 48u))}};
- return tint_symbol_13;
+ tint_padded_array_element arr[4] = (tint_padded_array_element[4])0;
+ {
+ for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
+ arr[i_1].el = tint_symbol_10(buffer, (offset + (i_1 * 16u)));
+ }
+ }
+ return arr;
}
[numthreads(1, 1, 1)]
@@ -48,7 +53,7 @@
const int2 g = asint(s[3].xy);
const int2 h = asint(s[3].zw);
const float2x3 i = tint_symbol_7(s, 64u);
- const float3x2 j = tint_symbol_9(s, 96u);
+ const float3x2 j = tint_symbol_8(s, 96u);
const Inner k = tint_symbol_10(s, 128u);
const tint_padded_array_element l[4] = tint_symbol_11(s, 144u);
return;
diff --git a/test/bug/tint/403.wgsl.expected.hlsl b/test/bug/tint/403.wgsl.expected.hlsl
index ef0848c..f798195 100644
--- a/test/bug/tint/403.wgsl.expected.hlsl
+++ b/test/bug/tint/403.wgsl.expected.hlsl
@@ -12,7 +12,7 @@
float4 value : SV_Position;
};
-float2x2 tint_symbol_4(uint4 buffer[1], uint offset) {
+float2x2 tint_symbol_3(uint4 buffer[1], uint offset) {
const uint scalar_offset = ((offset + 0u)) / 4;
uint4 ubo_load = buffer[scalar_offset / 4];
const uint scalar_offset_1 = ((offset + 8u)) / 4;
@@ -20,7 +20,7 @@
return float2x2(asfloat(((scalar_offset & 2) ? ubo_load.zw : ubo_load.xy)), asfloat(((scalar_offset_1 & 2) ? ubo_load_1.zw : ubo_load_1.xy)));
}
-float2x2 tint_symbol_6(uint4 buffer[1], uint offset) {
+float2x2 tint_symbol_5(uint4 buffer[1], uint offset) {
const uint scalar_offset_2 = ((offset + 0u)) / 4;
uint4 ubo_load_2 = buffer[scalar_offset_2 / 4];
const uint scalar_offset_3 = ((offset + 8u)) / 4;
@@ -31,8 +31,8 @@
tint_symbol_2 main(tint_symbol_1 tint_symbol) {
const uint gl_VertexIndex = tint_symbol.gl_VertexIndex;
float2 indexable[3] = (float2[3])0;
- const float2x2 x_23 = tint_symbol_4(x_20, 0u);
- const float2x2 x_28 = tint_symbol_6(x_26, 0u);
+ const float2x2 x_23 = tint_symbol_3(x_20, 0u);
+ const float2x2 x_28 = tint_symbol_5(x_26, 0u);
const uint x_46 = gl_VertexIndex;
const float2 tint_symbol_7[3] = {float2(-1.0f, 1.0f), float2(1.0f, 1.0f), float2(-1.0f, -1.0f)};
indexable = tint_symbol_7;
diff --git a/test/bug/tint/870.spvasm.expected.hlsl b/test/bug/tint/870.spvasm.expected.hlsl
index b3b5071..db19a18 100644
--- a/test/bug/tint/870.spvasm.expected.hlsl
+++ b/test/bug/tint/870.spvasm.expected.hlsl
@@ -1,14 +1,19 @@
ByteAddressBuffer sspp962805860buildInformation : register(t2, space0);
-typedef int tint_symbol_1_ret[6];
-tint_symbol_1_ret tint_symbol_1(ByteAddressBuffer buffer, uint offset) {
- const int tint_symbol_2[6] = {asint(buffer.Load((offset + 0u))), asint(buffer.Load((offset + 4u))), asint(buffer.Load((offset + 8u))), asint(buffer.Load((offset + 12u))), asint(buffer.Load((offset + 16u))), asint(buffer.Load((offset + 20u)))};
- return tint_symbol_2;
+typedef int tint_symbol_ret[6];
+tint_symbol_ret tint_symbol(ByteAddressBuffer buffer, uint offset) {
+ int arr[6] = (int[6])0;
+ {
+ for(uint i = 0u; (i < 6u); i = (i + 1u)) {
+ arr[i] = asint(buffer.Load((offset + (i * 4u))));
+ }
+ }
+ return arr;
}
void main_1() {
int orientation[6] = (int[6])0;
- const int x_23[6] = tint_symbol_1(sspp962805860buildInformation, 36u);
+ const int x_23[6] = tint_symbol(sspp962805860buildInformation, 36u);
orientation[0] = x_23[0u];
orientation[1] = x_23[1u];
orientation[2] = x_23[2u];
diff --git a/test/bug/tint/922.wgsl.expected.hlsl b/test/bug/tint/922.wgsl.expected.hlsl
index e0b2049..c3b7973 100644
--- a/test/bug/tint/922.wgsl.expected.hlsl
+++ b/test/bug/tint/922.wgsl.expected.hlsl
@@ -154,7 +154,7 @@
return o4;
}
-Mat4x3_ tint_symbol_5(uint4 buffer[96], uint offset) {
+Mat4x3_ tint_symbol_4(uint4 buffer[96], uint offset) {
const uint scalar_offset = ((offset + 0u)) / 4;
const uint scalar_offset_1 = ((offset + 16u)) / 4;
const uint scalar_offset_2 = ((offset + 32u)) / 4;
@@ -162,7 +162,7 @@
return tint_symbol_10;
}
-Mat4x4_ tint_symbol_7(uint4 buffer[4], uint offset) {
+Mat4x4_ tint_symbol_6(uint4 buffer[4], uint offset) {
const uint scalar_offset_3 = ((offset + 0u)) / 4;
const uint scalar_offset_4 = ((offset + 16u)) / 4;
const uint scalar_offset_5 = ((offset + 32u)) / 4;
@@ -181,13 +181,13 @@
void main1() {
Mat4x3_ t_PosMtx = (Mat4x3_)0;
float2 t_TexSpaceCoord = float2(0.0f, 0.0f);
- const Mat4x3_ _e18 = tint_symbol_5(global2, (48u * uint(int(a_PosMtxIdx1))));
+ const Mat4x3_ _e18 = tint_symbol_4(global2, (48u * uint(int(a_PosMtxIdx1))));
t_PosMtx = _e18;
const Mat4x4_ _e24 = _Mat4x4_1(t_PosMtx);
const float3 _e25 = a_Position1;
const Mat4x4_ _e30 = _Mat4x4_1(t_PosMtx);
const float4 _e34 = Mul(_e30, float4(a_Position1, 1.0f));
- const Mat4x4_ _e35 = tint_symbol_7(global, 0u);
+ const Mat4x4_ _e35 = tint_symbol_6(global, 0u);
const Mat4x4_ _e38 = _Mat4x4_1(t_PosMtx);
const float3 _e39 = a_Position1;
const Mat4x4_ _e44 = _Mat4x4_1(t_PosMtx);
diff --git a/test/bug/tint/998.wgsl.expected.hlsl b/test/bug/tint/998.wgsl.expected.hlsl
index 1acfe64..88f4d37 100644
--- a/test/bug/tint/998.wgsl.expected.hlsl
+++ b/test/bug/tint/998.wgsl.expected.hlsl
@@ -1,5 +1,3 @@
-SKIP: FAILED
-
cbuffer cbuffer_constants : register(b0, space1) {
uint4 constants[1];
};
@@ -17,5 +15,3 @@
s.data[constants[0].x] = 0u;
return;
}
-C:\src\tint\test\Shader@0x0000015D0E1BAC50(15,3-24): error X3500: array reference cannot be used as an l-value; not natively addressable
-
diff --git a/test/samples/cube.wgsl.expected.hlsl b/test/samples/cube.wgsl.expected.hlsl
index 9791b42..f83fd1f 100644
--- a/test/samples/cube.wgsl.expected.hlsl
+++ b/test/samples/cube.wgsl.expected.hlsl
@@ -19,7 +19,7 @@
float4 Position : SV_Position;
};
-float4x4 tint_symbol_7(uint4 buffer[4], uint offset) {
+float4x4 tint_symbol_6(uint4 buffer[4], uint offset) {
const uint scalar_offset = ((offset + 0u)) / 4;
const uint scalar_offset_1 = ((offset + 16u)) / 4;
const uint scalar_offset_2 = ((offset + 32u)) / 4;
@@ -30,7 +30,7 @@
tint_symbol_2 vtx_main(tint_symbol_1 tint_symbol) {
const VertexInput input = {tint_symbol.cur_position, tint_symbol.color};
VertexOutput output = (VertexOutput)0;
- output.Position = mul(input.cur_position, tint_symbol_7(uniforms, 0u));
+ output.Position = mul(input.cur_position, tint_symbol_6(uniforms, 0u));
output.vtxFragColor = input.color;
const tint_symbol_2 tint_symbol_8 = {output.vtxFragColor, output.Position};
return tint_symbol_8;
diff --git a/test/shader_io/shared_struct_storage_buffer.wgsl.expected.hlsl b/test/shader_io/shared_struct_storage_buffer.wgsl.expected.hlsl
index af4440e..c06b40f 100644
--- a/test/shader_io/shared_struct_storage_buffer.wgsl.expected.hlsl
+++ b/test/shader_io/shared_struct_storage_buffer.wgsl.expected.hlsl
@@ -12,7 +12,7 @@
float4 v : SV_Position;
};
-void tint_symbol_5(RWByteAddressBuffer buffer, uint offset, S value) {
+void tint_symbol_2(RWByteAddressBuffer buffer, uint offset, S value) {
buffer.Store((offset + 0u), asuint(value.f));
buffer.Store((offset + 4u), asuint(value.u));
buffer.Store4((offset + 128u), asuint(value.v));
@@ -23,6 +23,6 @@
const float f = input.f;
const uint u = input.u;
const float4 v = input.v;
- tint_symbol_5(output, 0u, input);
+ tint_symbol_2(output, 0u, input);
return;
}
