From 883fb63e0143317e74fa7087eb40b238ef6ab6a4 Mon Sep 17 00:00:00 2001 From: Ben Clayton Date: Fri, 16 Jul 2021 19:47:44 +0000 Subject: [PATCH] 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 Commit-Queue: Ben Clayton Reviewed-by: David Neto Auto-Submit: Ben Clayton --- src/transform/decompose_memory_access.cc | 259 ++++++++++-------- src/transform/decompose_memory_access_test.cc | 244 +++++++++-------- .../generator_impl_member_accessor_test.cc | 4 +- .../assign_to_function_var.wgsl.expected.hlsl | 39 +-- .../assign_to_private_var.wgsl.expected.hlsl | 39 +-- .../assign_to_storage_var.wgsl.expected.hlsl | 100 ++++--- ...assign_to_workgroup_var.wgsl.expected.hlsl | 39 +-- .../dynamic_index/read.wgsl.expected.hlsl | 19 +- .../dynamic_index/write.wgsl.expected.hlsl | 18 +- .../static_index/read.wgsl.expected.hlsl | 13 +- .../static_index/write.wgsl.expected.hlsl | 14 +- .../static_index/read.wgsl.expected.hlsl | 13 +- test/bug/tint/403.wgsl.expected.hlsl | 8 +- test/bug/tint/870.spvasm.expected.hlsl | 15 +- test/bug/tint/922.wgsl.expected.hlsl | 8 +- test/bug/tint/998.wgsl.expected.hlsl | 4 - test/samples/cube.wgsl.expected.hlsl | 4 +- ...d_struct_storage_buffer.wgsl.expected.hlsl | 4 +- 18 files changed, 476 insertions(+), 368 deletions(-) diff --git a/src/transform/decompose_memory_access.cc b/src/transform/decompose_memory_access.cc index 06251dedf1..f06b3ae168 100644 --- a/src/transform/decompose_memory_access.cc +++ b/src/transform/decompose_memory_access.cc @@ -303,6 +303,10 @@ struct Store { /// 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 @@ struct DecomposeMemoryAccess::State { /// Allocations for offsets BlockAllocator 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 @@ struct DecomposeMemoryAccess::State { /// 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 @@ struct DecomposeMemoryAccess::State { load_funcs, LoadStoreKey{storage_class, buf_ty, el_ty}, [&] { auto* buf_ast_ty = CreateASTTypeFor(ctx, buf_ty); auto* disable_validation = - ctx.dst->ASTNodes().Create( - ctx.dst->ID(), ast::DisabledValidation:: - kIgnoreConstructibleFunctionParameter); + b.ASTNodes().Create( + 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( - 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(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( - ctx.dst->Sym(), params, el_ast_ty, nullptr, + auto* func = b.create( + name, params, el_ast_ty, nullptr, ast::DecorationList{ intrinsic, - ctx.dst->ASTNodes() - .Create( - ctx.dst->ID(), - ast::DisabledValidation::kFunctionHasNoBody), + b.ASTNodes().Create( + b.ID(), ast::DisabledValidation::kFunctionHasNoBody), }, ast::DecorationList{}); + b.AST().AddFunction(func); + } else if (auto* arr_ty = el_ty->As()) { + // fn load_func(buf : buf_ty, offset : u32) -> array { + // var arr : array; + // 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::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()) { 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()) { 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()) { - 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( - ctx.dst->Sym(), params, el_ast_ty, - ctx.dst->Block(ctx.dst->Return( - ctx.dst->create( - CreateASTTypeFor(ctx, el_ty), values))), - ast::DecorationList{}, ast::DecorationList{}); + b.Func(name, params, CreateASTTypeFor(ctx, el_ty), + { + b.Return(b.create( + CreateASTTypeFor(ctx, el_ty), values)), + }); } - ctx.dst->AST().AddFunction(func); - return func->symbol(); + return name; }); } @@ -525,13 +550,11 @@ struct DecomposeMemoryAccess::State { /// 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 @@ struct DecomposeMemoryAccess::State { auto* buf_ast_ty = CreateASTTypeFor(ctx, buf_ty); auto* el_ast_ty = CreateASTTypeFor(ctx, el_ty); auto* disable_validation = - ctx.dst->ASTNodes().Create( - ctx.dst->ID(), ast::DisabledValidation:: - kIgnoreConstructibleFunctionParameter); + b.ASTNodes().Create( + 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( - 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(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( - ctx.dst->Sym(), params, ctx.dst->ty.void_(), nullptr, + auto* func = b.create( + name, params, b.ty.void_(), nullptr, ast::DecorationList{ intrinsic, - ctx.dst->ASTNodes() - .Create( - ctx.dst->ID(), - ast::DisabledValidation::kFunctionHasNoBody), + b.ASTNodes().Create( + b.ID(), ast::DisabledValidation::kFunctionHasNoBody), }, ast::DecorationList{}); - + b.AST().AddFunction(func); } else { ast::StatementList body; - if (auto* mat_ty = el_ty->As()) { + if (auto* arr_ty = el_ty->As()) { + // 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::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( + 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()) { 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(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(call)); } } else if (auto* str = el_ty->As()) { 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(call)); - } - } else if (auto* arr = el_ty->As()) { - 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(call)); + Symbol store = + StoreFunc(buf_ty, member->Type()->UnwrapRef(), var_user); + auto* call = b.Call(store, "buffer", offset, access); + body.emplace_back(b.create(call)); } } - func = ctx.dst->create( - 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 @@ struct DecomposeMemoryAccess::State { /// 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 @@ struct DecomposeMemoryAccess::State { 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( - ctx.dst->ID(), + b.ASTNodes().Create( + 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 @@ struct DecomposeMemoryAccess::State { ast::VariableList params = { // Note: The buffer parameter requires the kStorage StorageClass in // order for HLSL to emit this as a ByteAddressBuffer. - ctx.dst->create( - 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(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( - ctx.dst->Sym(), params, ret_ty, nullptr, + auto* func = b.create( + b.Sym(), params, ret_ty, nullptr, ast::DecorationList{ atomic, - ctx.dst->ASTNodes().Create( - ctx.dst->ID(), ast::DisabledValidation::kFunctionHasNoBody), + b.ASTNodes().Create( + b.ID(), ast::DisabledValidation::kFunctionHasNoBody), }, ast::DecorationList{}); - ctx.dst->AST().AddFunction(func); + b.AST().AddFunction(func); return func->symbol(); }); } @@ -777,7 +810,7 @@ DecomposeMemoryAccess::~DecomposeMemoryAccess() = default; 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 @@ void DecomposeMemoryAccess::Run(CloneContext& ctx, const DataMap&, DataMap&) { auto* buf_ty = access.var->Type()->UnwrapRef(); auto* el_ty = access.type->UnwrapRef()->As()->Type(); Symbol func = - state.AtomicFunc(ctx, buf_ty, el_ty, intrinsic, + state.AtomicFunc(buf_ty, el_ty, intrinsic, access.var->As()); ast::ExpressionList args{ctx.Clone(buf), offset}; @@ -937,8 +970,8 @@ void DecomposeMemoryAccess::Run(CloneContext& ctx, const DataMap&, DataMap&) { 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()); + Symbol func = + state.LoadFunc(buf_ty, el_ty, access.var->As()); return ctx.dst->Call(func, ctx.CloneWithoutTransform(buf), offset); }); } @@ -951,7 +984,7 @@ void DecomposeMemoryAccess::Run(CloneContext& ctx, const DataMap&, DataMap&) { 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()); 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 5e9e60da2e..58466c4e94 100644 --- a/src/transform/decompose_memory_access_test.cc +++ b/src/transform/decompose_memory_access_test.cc @@ -181,7 +181,11 @@ fn tint_symbol_20([[internal(disable_validation__ignore_constructible_function_p } fn tint_symbol_21([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> array, 2> { - return array, 2>(tint_symbol_8(buffer, (offset + 0u)), tint_symbol_8(buffer, (offset + 16u))); + var arr : array, 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_20([[internal(disable_validation__ignore_constructible_function_p } fn tint_symbol_21([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : UB, offset : u32) -> array, 2> { - return array, 2>(tint_symbol_8(buffer, (offset + 0u)), tint_symbol_8(buffer, (offset + 16u))); + var arr : array, 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_20([[internal(disable_validation__ignore_constructible_function_p } fn tint_symbol_21([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : array, 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 @@ struct SB { [[group(0), binding(0)]] var 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 +fn tint_symbol_4([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec2 [[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 +fn tint_symbol_5([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec2 [[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 +fn tint_symbol_6([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec2 [[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 +fn tint_symbol_7([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec3 [[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 +fn tint_symbol_8([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec3 [[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 +fn tint_symbol_9([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec3 [[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 +fn tint_symbol_10([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec4 [[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 +fn tint_symbol_11([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec4 [[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 +fn tint_symbol_12([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec4 -fn tint_symbol_12([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat2x2 { - return mat2x2(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 { + return mat2x2(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 { - return mat2x3(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 { + return mat2x3(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 { - return mat2x4(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 { + return mat2x4(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 { - return mat3x2(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 { + return mat3x2(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 { - return mat3x3(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 { + return mat3x3(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 { - return mat3x4(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 { + return mat3x4(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 { - return mat4x2(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 { + return mat4x2(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 { - return mat4x3(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 { + return mat4x3(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 { - return mat4x4(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 { + return mat4x4(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, 2> { - return array, 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, 2> { + var arr : array, 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 @@ struct SB { [[group(0), binding(0)]] var 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) +fn tint_symbol_4([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec2) [[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) +fn tint_symbol_5([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec2) [[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) +fn tint_symbol_6([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec2) [[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) +fn tint_symbol_7([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec3) [[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) +fn tint_symbol_8([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec3) [[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) +fn tint_symbol_9([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec3) [[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) +fn tint_symbol_10([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec4) [[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) +fn tint_symbol_11([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec4) [[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) +fn tint_symbol_12([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec4) -fn tint_symbol_12([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat2x2) { - 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) { + 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) { - 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) { + 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) { - 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) { + 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) { - 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) { + 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) { - 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) { + 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) { - 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) { + 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) { - 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) { + 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) { - 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) { + 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) { - 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) { + 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, 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, 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 10d7cb91d0..40b73d4259 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 @@ TEST_F(HlslGeneratorImplTest_MemberAccessor, StorageBuffer_Store_Matrix_Empty) { 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 2c7d6be159..5b15fc1256 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 @@ S ret_struct_arr() { 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 @@ void foo(tint_padded_array_element src_param[4]) { 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 68ea5c3f47..8cd8ee1b62 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 @@ S ret_struct_arr() { 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 @@ void foo(tint_padded_array_element src_param[4]) { 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 ee58066595..8825ef39d7 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 @@ S ret_struct_arr() { 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 48c5fae141..43a8a1cb29 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 @@ S ret_struct_arr() { 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 @@ void foo(tint_padded_array_element src_param[4]) { 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 9b9390af18..24312d1f54 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 @@ float2x3 tint_symbol_8(ByteAddressBuffer buffer, uint offset) { 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 @@ void main(tint_symbol_1 tint_symbol) { 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 74aeaf2068..c18b7c6845 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 @@ void tint_symbol_8(RWByteAddressBuffer buffer, uint offset, float2x3 value) { 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 @@ void main(tint_symbol_1 tint_symbol) { 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 c14caff8f4..42a6a71d42 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 @@ float2x3 tint_symbol_6(ByteAddressBuffer buffer, uint offset) { 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 @@ Inner tint_symbol_9(ByteAddressBuffer buffer, uint offset) { 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 @@ void main() { 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 7e569cdaa2..8c9b7ea6c5 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 @@ void tint_symbol_6(RWByteAddressBuffer buffer, uint offset, float2x3 value) { 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_9(RWByteAddressBuffer buffer, uint offset, Inner value) { } 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 @@ void main() { 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 fa6d133619..f66c81b288 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 @@ float2x3 tint_symbol_7(uint4 buffer[13], uint offset) { 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 @@ Inner tint_symbol_10(uint4 buffer[13], uint offset) { 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 @@ void main() { 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 ef0848ccce..f7981958e9 100644 --- a/test/bug/tint/403.wgsl.expected.hlsl +++ b/test/bug/tint/403.wgsl.expected.hlsl @@ -12,7 +12,7 @@ struct tint_symbol_2 { 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 @@ float2x2 tint_symbol_4(uint4 buffer[1], uint offset) { 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 @@ float2x2 tint_symbol_6(uint4 buffer[1], uint offset) { 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 b3b5071b52..db19a18586 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 e0b2049f07..c3b7973e88 100644 --- a/test/bug/tint/922.wgsl.expected.hlsl +++ b/test/bug/tint/922.wgsl.expected.hlsl @@ -154,7 +154,7 @@ Mat4x3_ _Mat4x3_1(Mat4x4_ m20) { 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 @@ Mat4x3_ tint_symbol_5(uint4 buffer[96], uint offset) { 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 @@ Mat4x2_ tint_symbol_9(uint4 buffer[3], uint offset) { 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 1acfe64bee..88f4d37b64 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 @@ void main() { 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 9791b42083..f83fd1f5e0 100644 --- a/test/samples/cube.wgsl.expected.hlsl +++ b/test/samples/cube.wgsl.expected.hlsl @@ -19,7 +19,7 @@ struct tint_symbol_2 { 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 @@ float4x4 tint_symbol_7(uint4 buffer[4], uint offset) { 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 af4440eae8..c06b40f372 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 @@ struct tint_symbol_1 { 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 @@ void frag_main(tint_symbol_1 tint_symbol) { 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; }