From 2032d034002a8002774a996b4707ca6b082949c6 Mon Sep 17 00:00:00 2001 From: Ben Clayton Date: Wed, 15 Jun 2022 19:32:37 +0000 Subject: [PATCH] tint/transform: Remove use of StorageClass on parameter Parameters don't have storage classes or access qualifiers. This was just (ab)using the fact that a parameter uses the same AST type as a 'var'. Also simplify the parameter disable validation logic. Bug: tint:1582 Change-Id: Ic218078a410f991e7956e6cb23621a94a69b75a3 Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/93603 Commit-Queue: Ben Clayton Reviewed-by: Dan Sinclair --- src/tint/ast/disable_validation_attribute.cc | 4 +- src/tint/ast/disable_validation_attribute.h | 28 +- src/tint/resolver/validator.cc | 26 +- src/tint/transform/calculate_array_length.cc | 58 +- .../transform/calculate_array_length_test.cc | 60 +- src/tint/transform/decompose_memory_access.cc | 337 +++---- .../transform/decompose_memory_access_test.cc | 920 +++++++++--------- src/tint/transform/manager.cc | 3 +- src/tint/writer/hlsl/generator_impl.cc | 23 +- 9 files changed, 734 insertions(+), 725 deletions(-) diff --git a/src/tint/ast/disable_validation_attribute.cc b/src/tint/ast/disable_validation_attribute.cc index 4bc9f74c65..c5a6545b68 100644 --- a/src/tint/ast/disable_validation_attribute.cc +++ b/src/tint/ast/disable_validation_attribute.cc @@ -35,8 +35,8 @@ std::string DisableValidationAttribute::InternalName() const { return "disable_validation__ignore_storage_class"; case DisabledValidation::kEntryPointParameter: return "disable_validation__entry_point_parameter"; - case DisabledValidation::kIgnoreConstructibleFunctionParameter: - return "disable_validation__ignore_constructible_function_parameter"; + case DisabledValidation::kFunctionParameter: + return "disable_validation__function_parameter"; case DisabledValidation::kIgnoreStrideAttribute: return "disable_validation__ignore_stride"; case DisabledValidation::kIgnoreInvalidPointerArgument: diff --git a/src/tint/ast/disable_validation_attribute.h b/src/tint/ast/disable_validation_attribute.h index db70ad4c9a..e44f7b8906 100644 --- a/src/tint/ast/disable_validation_attribute.h +++ b/src/tint/ast/disable_validation_attribute.h @@ -24,27 +24,23 @@ namespace tint::ast { /// Enumerator of validation features that can be disabled with a /// DisableValidationAttribute attribute. enum class DisabledValidation { - /// When applied to a function, the validator will not complain there is no - /// body to a function. + /// When applied to a function, the validator will not complain there is no body to a function. kFunctionHasNoBody, - /// When applied to a module-scoped variable, the validator will not complain - /// if two resource variables have the same binding points. + /// When applied to a module-scoped variable, the validator will not complain if two resource + /// variables have the same binding points. kBindingPointCollision, - /// When applied to a variable, the validator will not complain about the - /// declared storage class. + /// When applied to a variable, the validator will not complain about the declared storage + /// class. kIgnoreStorageClass, - /// When applied to an entry-point function parameter, the validator will not - /// check for entry IO attributes. + /// When applied to an entry-point function parameter, the validator will not check for entry IO + /// attributes. kEntryPointParameter, - /// When applied to a function parameter, the validator will not - /// check if parameter type is constructible - kIgnoreConstructibleFunctionParameter, - /// When applied to a member attribute, a stride attribute may be applied to - /// non-array types. + /// When applied to a function parameter, the parameter will not be validated. + kFunctionParameter, + /// When applied to a member attribute, a stride attribute may be applied to non-array types. kIgnoreStrideAttribute, - /// When applied to a pointer function parameter, the validator will not - /// require a function call argument passed for that parameter to have a - /// certain form. + /// When applied to a pointer function parameter, the validator will not require a function call + /// argument passed for that parameter to have a certain form. kIgnoreInvalidPointerArgument, }; diff --git a/src/tint/resolver/validator.cc b/src/tint/resolver/validator.cc index 2b4d8a99d3..4d1ee075b7 100644 --- a/src/tint/resolver/validator.cc +++ b/src/tint/resolver/validator.cc @@ -722,19 +722,20 @@ bool Validator::FunctionParameter(const ast::Function* func, const sem::Variable auto* decl = var->Declaration(); + if (IsValidationDisabled(decl->attributes, ast::DisabledValidation::kFunctionParameter)) { + return true; + } + for (auto* attr : decl->attributes) { if (!func->IsEntryPoint() && !attr->Is()) { AddError("attribute is not valid for non-entry point function parameters", attr->source); return false; - } else if (!attr->IsAnyOf() && - (IsValidationEnabled(decl->attributes, - ast::DisabledValidation::kEntryPointParameter) && - IsValidationEnabled( - decl->attributes, - ast::DisabledValidation::kIgnoreConstructibleFunctionParameter))) { + } + if (!attr->IsAnyOf() && + (IsValidationEnabled(decl->attributes, + ast::DisabledValidation::kEntryPointParameter))) { AddError("attribute is not valid for function parameters", attr->source); return false; } @@ -753,9 +754,7 @@ bool Validator::FunctionParameter(const ast::Function* func, const sem::Variable } if (IsPlain(var->Type())) { - if (!var->Type()->IsConstructible() && - IsValidationEnabled(decl->attributes, - ast::DisabledValidation::kIgnoreConstructibleFunctionParameter)) { + if (!var->Type()->IsConstructible()) { AddError("store type of function parameter must be a constructible type", decl->source); return false; } @@ -964,9 +963,8 @@ bool Validator::Function(const sem::Function* func, ast::PipelineStage stage) co ast::InvariantAttribute>() && (IsValidationEnabled(decl->attributes, ast::DisabledValidation::kEntryPointParameter) && - IsValidationEnabled( - decl->attributes, - ast::DisabledValidation::kIgnoreConstructibleFunctionParameter))) { + IsValidationEnabled(decl->attributes, + ast::DisabledValidation::kFunctionParameter))) { AddError("attribute is not valid for entry point return types", attr->source); return false; } diff --git a/src/tint/transform/calculate_array_length.cc b/src/tint/transform/calculate_array_length.cc index 9cf1175970..acf55a688a 100644 --- a/src/tint/transform/calculate_array_length.cc +++ b/src/tint/transform/calculate_array_length.cc @@ -23,6 +23,7 @@ #include "src/tint/sem/block_statement.h" #include "src/tint/sem/call.h" #include "src/tint/sem/function.h" +#include "src/tint/sem/reference.h" #include "src/tint/sem/statement.h" #include "src/tint/sem/struct.h" #include "src/tint/sem/variable.h" @@ -89,22 +90,20 @@ void CalculateArrayLength::Run(CloneContext& ctx, const DataMap&, DataMap&) cons // get_buffer_size_intrinsic() emits the function decorated with // BufferSizeIntrinsic that is transformed by the HLSL writer into a call to // [RW]ByteAddressBuffer.GetDimensions(). - std::unordered_map buffer_size_intrinsics; - auto get_buffer_size_intrinsic = [&](const sem::Type* buffer_type) { + std::unordered_map buffer_size_intrinsics; + auto get_buffer_size_intrinsic = [&](const sem::Reference* buffer_type) { return utils::GetOrCreate(buffer_size_intrinsics, buffer_type, [&] { auto name = ctx.dst->Sym(); auto* type = CreateASTTypeFor(ctx, buffer_type); auto* disable_validation = - ctx.dst->Disable(ast::DisabledValidation::kIgnoreConstructibleFunctionParameter); + ctx.dst->Disable(ast::DisabledValidation::kFunctionParameter); ctx.dst->AST().AddFunction(ctx.dst->create( name, ast::ParameterList{ - // 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, - ast::Access::kUndefined, type, true, false, - nullptr, ast::AttributeList{disable_validation}), + ctx.dst->Param("buffer", + ctx.dst->ty.pointer(type, buffer_type->StorageClass(), + buffer_type->Access()), + {disable_validation}), ctx.dst->Param("result", ctx.dst->ty.pointer(ctx.dst->ty.u32(), ast::StorageClass::kFunction)), }, @@ -128,10 +127,10 @@ void CalculateArrayLength::Run(CloneContext& ctx, const DataMap&, DataMap&) cons if (builtin->Type() == sem::BuiltinType::kArrayLength) { // We're dealing with an arrayLength() call - // A runtime-sized array can only appear as the store type of a - // variable, or the last element of a structure (which cannot itself - // be nested). Given that we require SimplifyPointers, we can assume - // that the arrayLength() call has one of two forms: + // A runtime-sized array can only appear as the store type of a variable, or the + // last element of a structure (which cannot itself be nested). Given that we + // require SimplifyPointers, we can assume that the arrayLength() call has one + // of two forms: // arrayLength(&struct_var.array_member) // arrayLength(&array_var) auto* arg = call_expr->args[0]; @@ -152,10 +151,9 @@ void CalculateArrayLength::Run(CloneContext& ctx, const DataMap&, DataMap&) cons break; } auto* storage_buffer_var = storage_buffer_sem->Variable(); - auto* storage_buffer_type = storage_buffer_sem->Type()->UnwrapRef(); + auto* storage_buffer_type = storage_buffer_sem->Type()->As(); - // Generate BufferSizeIntrinsic for this storage type if we haven't - // already + // Generate BufferSizeIntrinsic for this storage type if we haven't already auto buffer_size = get_buffer_size_intrinsic(storage_buffer_type); // Find the current statement block @@ -177,7 +175,7 @@ void CalculateArrayLength::Run(CloneContext& ctx, const DataMap&, DataMap&) cons // BufferSizeIntrinsic(X, ARGS...) is // translated to: // X.GetDimensions(ARGS..) by the writer - buffer_size, ctx.Clone(storage_buffer_expr), + buffer_size, ctx.dst->AddressOf(ctx.Clone(storage_buffer_expr)), ctx.dst->AddressOf( ctx.dst->Expr(buffer_size_result->variable->symbol)))); @@ -188,22 +186,26 @@ void CalculateArrayLength::Run(CloneContext& ctx, const DataMap&, DataMap&) cons auto name = ctx.dst->Sym(); const ast::Expression* total_size = ctx.dst->Expr(buffer_size_result->variable); - const sem::Array* array_type = nullptr; - if (auto* str = storage_buffer_type->As()) { - // The variable is a struct, so subtract the byte offset of - // the array member. - auto* array_member_sem = str->Members().back(); - array_type = array_member_sem->Type()->As(); - total_size = - ctx.dst->Sub(total_size, u32(array_member_sem->Offset())); - } else if (auto* arr = storage_buffer_type->As()) { - array_type = arr; - } else { + + const sem::Array* array_type = Switch( + storage_buffer_type->StoreType(), + [&](const sem::Struct* str) { + // The variable is a struct, so subtract the byte offset of + // the array member. + auto* array_member_sem = str->Members().back(); + total_size = + ctx.dst->Sub(total_size, u32(array_member_sem->Offset())); + return array_member_sem->Type()->As(); + }, + [&](const sem::Array* arr) { return arr; }); + + if (!array_type) { TINT_ICE(Transform, ctx.dst->Diagnostics()) << "expected form of arrayLength argument to be " "&array_var or &struct_var.array_member"; return name; } + uint32_t array_stride = array_type->Size(); auto* array_length_var = ctx.dst->Decl( ctx.dst->Let(name, ctx.dst->ty.u32(), diff --git a/src/tint/transform/calculate_array_length_test.cc b/src/tint/transform/calculate_array_length_test.cc index e2674b0572..98f2ed753d 100644 --- a/src/tint/transform/calculate_array_length_test.cc +++ b/src/tint/transform/calculate_array_length_test.cc @@ -76,14 +76,14 @@ fn main() { auto* expect = R"( @internal(intrinsic_buffer_size) -fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : array, result : ptr) +fn tint_symbol(@internal(disable_validation__function_parameter) buffer : ptr, read>, result : ptr) @group(0) @binding(0) var sb : array; @compute @workgroup_size(1) fn main() { var tint_symbol_1 : u32 = 0u; - tint_symbol(sb, &(tint_symbol_1)); + tint_symbol(&(sb), &(tint_symbol_1)); let tint_symbol_2 : u32 = (tint_symbol_1 / 4u); var len : u32 = tint_symbol_2; } @@ -111,7 +111,7 @@ fn main() { auto* expect = R"( @internal(intrinsic_buffer_size) -fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, result : ptr) +fn tint_symbol(@internal(disable_validation__function_parameter) buffer : ptr, result : ptr) struct SB { x : i32, @@ -123,7 +123,7 @@ struct SB { @compute @workgroup_size(1) fn main() { var tint_symbol_1 : u32 = 0u; - tint_symbol(sb, &(tint_symbol_1)); + tint_symbol(&(sb), &(tint_symbol_1)); let tint_symbol_2 : u32 = ((tint_symbol_1 - 4u) / 4u); var len : u32 = tint_symbol_2; } @@ -149,7 +149,7 @@ fn main() { )"; auto* expect = R"( @internal(intrinsic_buffer_size) -fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : array, result : ptr) +fn tint_symbol(@internal(disable_validation__function_parameter) buffer : ptr, read>, result : ptr) struct S { f : f32, @@ -160,7 +160,7 @@ struct S { @compute @workgroup_size(1) fn main() { var tint_symbol_1 : u32 = 0u; - tint_symbol(arr, &(tint_symbol_1)); + tint_symbol(&(arr), &(tint_symbol_1)); let tint_symbol_2 : u32 = (tint_symbol_1 / 4u); let len = tint_symbol_2; } @@ -186,7 +186,7 @@ fn main() { )"; auto* expect = R"( @internal(intrinsic_buffer_size) -fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : array>, result : ptr) +fn tint_symbol(@internal(disable_validation__function_parameter) buffer : ptr>, read>, result : ptr) struct S { f : f32, @@ -197,7 +197,7 @@ struct S { @compute @workgroup_size(1) fn main() { var tint_symbol_1 : u32 = 0u; - tint_symbol(arr, &(tint_symbol_1)); + tint_symbol(&(arr), &(tint_symbol_1)); let tint_symbol_2 : u32 = (tint_symbol_1 / 16u); let len = tint_symbol_2; } @@ -222,14 +222,14 @@ fn main() { auto* expect = R"( @internal(intrinsic_buffer_size) -fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : array, result : ptr) +fn tint_symbol(@internal(disable_validation__function_parameter) buffer : ptr, read>, result : ptr) @group(0) @binding(0) var sb : array; @compute @workgroup_size(1) fn main() { var tint_symbol_1 : u32 = 0u; - tint_symbol(sb, &(tint_symbol_1)); + tint_symbol(&(sb), &(tint_symbol_1)); let tint_symbol_2 : u32 = (tint_symbol_1 / 4u); var a : u32 = tint_symbol_2; var b : u32 = tint_symbol_2; @@ -261,7 +261,7 @@ fn main() { auto* expect = R"( @internal(intrinsic_buffer_size) -fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, result : ptr) +fn tint_symbol(@internal(disable_validation__function_parameter) buffer : ptr, result : ptr) struct SB { x : i32, @@ -273,7 +273,7 @@ struct SB { @compute @workgroup_size(1) fn main() { var tint_symbol_1 : u32 = 0u; - tint_symbol(sb, &(tint_symbol_1)); + tint_symbol(&(sb), &(tint_symbol_1)); let tint_symbol_2 : u32 = ((tint_symbol_1 - 4u) / 4u); var a : u32 = tint_symbol_2; var b : u32 = tint_symbol_2; @@ -309,7 +309,7 @@ fn main() { auto* expect = R"( @internal(intrinsic_buffer_size) -fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, result : ptr) +fn tint_symbol(@internal(disable_validation__function_parameter) buffer : ptr, result : ptr) struct SB { x : i32, @@ -322,13 +322,13 @@ struct SB { fn main() { if (true) { var tint_symbol_1 : u32 = 0u; - tint_symbol(sb, &(tint_symbol_1)); + tint_symbol(&(sb), &(tint_symbol_1)); let tint_symbol_2 : u32 = ((tint_symbol_1 - 4u) / 4u); var len : u32 = tint_symbol_2; } else { if (true) { var tint_symbol_3 : u32 = 0u; - tint_symbol(sb, &(tint_symbol_3)); + tint_symbol(&(sb), &(tint_symbol_3)); let tint_symbol_4 : u32 = ((tint_symbol_3 - 4u) / 4u); var len : u32 = tint_symbol_4; } @@ -370,13 +370,13 @@ fn main() { auto* expect = R"( @internal(intrinsic_buffer_size) -fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB1, result : ptr) +fn tint_symbol(@internal(disable_validation__function_parameter) buffer : ptr, result : ptr) @internal(intrinsic_buffer_size) -fn tint_symbol_3(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB2, result : ptr) +fn tint_symbol_3(@internal(disable_validation__function_parameter) buffer : ptr, result : ptr) @internal(intrinsic_buffer_size) -fn tint_symbol_6(@internal(disable_validation__ignore_constructible_function_parameter) buffer : array, result : ptr) +fn tint_symbol_6(@internal(disable_validation__function_parameter) buffer : ptr, read>, result : ptr) struct SB1 { x : i32, @@ -397,13 +397,13 @@ struct SB2 { @compute @workgroup_size(1) fn main() { var tint_symbol_1 : u32 = 0u; - tint_symbol(sb1, &(tint_symbol_1)); + tint_symbol(&(sb1), &(tint_symbol_1)); let tint_symbol_2 : u32 = ((tint_symbol_1 - 4u) / 4u); var tint_symbol_4 : u32 = 0u; - tint_symbol_3(sb2, &(tint_symbol_4)); + tint_symbol_3(&(sb2), &(tint_symbol_4)); let tint_symbol_5 : u32 = ((tint_symbol_4 - 16u) / 16u); var tint_symbol_7 : u32 = 0u; - tint_symbol_6(sb3, &(tint_symbol_7)); + tint_symbol_6(&(sb3), &(tint_symbol_7)); let tint_symbol_8 : u32 = (tint_symbol_7 / 4u); var len1 : u32 = tint_symbol_2; var len2 : u32 = tint_symbol_5; @@ -440,7 +440,7 @@ fn main() { auto* expect = R"( @internal(intrinsic_buffer_size) -fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, result : ptr) +fn tint_symbol(@internal(disable_validation__function_parameter) buffer : ptr, result : ptr) struct SB { x : i32, @@ -454,12 +454,12 @@ struct SB { @compute @workgroup_size(1) fn main() { var tint_symbol_1 : u32 = 0u; - tint_symbol(a, &(tint_symbol_1)); + tint_symbol(&(a), &(tint_symbol_1)); let tint_symbol_2 : u32 = ((tint_symbol_1 - 4u) / 4u); var a_1 : u32 = tint_symbol_2; { var tint_symbol_3 : u32 = 0u; - tint_symbol(a, &(tint_symbol_3)); + tint_symbol(&(a), &(tint_symbol_3)); let tint_symbol_4 : u32 = ((tint_symbol_3 - 4u) / 4u); var b_1 : u32 = tint_symbol_4; } @@ -500,24 +500,24 @@ struct SB2 { auto* expect = R"( @internal(intrinsic_buffer_size) -fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB1, result : ptr) +fn tint_symbol(@internal(disable_validation__function_parameter) buffer : ptr, result : ptr) @internal(intrinsic_buffer_size) -fn tint_symbol_3(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB2, result : ptr) +fn tint_symbol_3(@internal(disable_validation__function_parameter) buffer : ptr, result : ptr) @internal(intrinsic_buffer_size) -fn tint_symbol_6(@internal(disable_validation__ignore_constructible_function_parameter) buffer : array, result : ptr) +fn tint_symbol_6(@internal(disable_validation__function_parameter) buffer : ptr, read>, result : ptr) @compute @workgroup_size(1) fn main() { var tint_symbol_1 : u32 = 0u; - tint_symbol(sb1, &(tint_symbol_1)); + tint_symbol(&(sb1), &(tint_symbol_1)); let tint_symbol_2 : u32 = ((tint_symbol_1 - 4u) / 4u); var tint_symbol_4 : u32 = 0u; - tint_symbol_3(sb2, &(tint_symbol_4)); + tint_symbol_3(&(sb2), &(tint_symbol_4)); let tint_symbol_5 : u32 = ((tint_symbol_4 - 16u) / 16u); var tint_symbol_7 : u32 = 0u; - tint_symbol_6(sb3, &(tint_symbol_7)); + tint_symbol_6(&(sb3), &(tint_symbol_7)); let tint_symbol_8 : u32 = (tint_symbol_7 / 4u); var len1 : u32 = tint_symbol_2; var len2 : u32 = tint_symbol_5; diff --git a/src/tint/transform/decompose_memory_access.cc b/src/tint/transform/decompose_memory_access.cc index f89d9abbdc..32ad153bbe 100644 --- a/src/tint/transform/decompose_memory_access.cc +++ b/src/tint/transform/decompose_memory_access.cc @@ -98,29 +98,32 @@ struct OffsetBinOp : Offset { /// LoadStoreKey is the unordered map key to a load or store intrinsic. struct LoadStoreKey { ast::StorageClass const storage_class; // buffer storage class + ast::Access const access; // buffer access sem::Type const* buf_ty = nullptr; // buffer type sem::Type const* el_ty = nullptr; // element type bool operator==(const LoadStoreKey& rhs) const { - return storage_class == rhs.storage_class && buf_ty == rhs.buf_ty && el_ty == rhs.el_ty; + return storage_class == rhs.storage_class && access == rhs.access && buf_ty == rhs.buf_ty && + el_ty == rhs.el_ty; } struct Hasher { inline std::size_t operator()(const LoadStoreKey& u) const { - return utils::Hash(u.storage_class, u.buf_ty, u.el_ty); + return utils::Hash(u.storage_class, u.access, u.buf_ty, u.el_ty); } }; }; /// AtomicKey is the unordered map key to an atomic intrinsic. struct AtomicKey { + ast::Access const access; // buffer access sem::Type const* buf_ty = nullptr; // buffer type sem::Type const* el_ty = nullptr; // element type sem::BuiltinType const op; // atomic op bool operator==(const AtomicKey& rhs) const { - return buf_ty == rhs.buf_ty && el_ty == rhs.el_ty && op == rhs.op; + return access == rhs.access && buf_ty == rhs.buf_ty && el_ty == rhs.el_ty && op == rhs.op; } struct Hasher { inline std::size_t operator()(const AtomicKey& u) const { - return utils::Hash(u.buf_ty, u.el_ty, u.op); + return utils::Hash(u.access, u.buf_ty, u.el_ty, u.op); } }; }; @@ -420,10 +423,10 @@ struct DecomposeMemoryAccess::State { return access; } - /// LoadFunc() returns a symbol to an intrinsic function that loads an element - /// of type `el_ty` from a storage or uniform buffer of type `buf_ty`. + /// LoadFunc() returns a symbol to an intrinsic function that loads an element 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` + /// `fn load(buf : ptr, offset : u32) -> el_ty` /// @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 @@ -432,89 +435,84 @@ struct DecomposeMemoryAccess::State { const sem::Type* el_ty, const sem::VariableUser* var_user) { auto storage_class = var_user->Variable()->StorageClass(); - return utils::GetOrCreate(load_funcs, LoadStoreKey{storage_class, buf_ty, el_ty}, [&] { - auto* buf_ast_ty = CreateASTTypeFor(ctx, buf_ty); - auto* disable_validation = - b.Disable(ast::DisabledValidation::kIgnoreConstructibleFunctionParameter); + auto access = var_user->Variable()->Access(); + return utils::GetOrCreate( + load_funcs, LoadStoreKey{storage_class, access, buf_ty, el_ty}, [&] { + ast::ParameterList params = { + b.Param("buffer", + b.ty.pointer(CreateASTTypeFor(ctx, buf_ty), storage_class, access), + {b.Disable(ast::DisabledValidation::kFunctionParameter)}), + b.Param("offset", b.ty.u32()), + }; - ast::ParameterList params = { - // Note: The buffer parameter requires the StorageClass in - // order for HLSL to emit this as a ByteAddressBuffer or cbuffer - // array. - b.create(b.Sym("buffer"), storage_class, - var_user->Variable()->Access(), buf_ast_ty, true, false, - nullptr, ast::AttributeList{disable_validation}), - b.Param("offset", b.ty.u32()), - }; + auto name = b.Sym(); - auto name = b.Sym(); + if (auto* intrinsic = IntrinsicLoadFor(ctx.dst, storage_class, el_ty)) { + auto* el_ast_ty = CreateASTTypeFor(ctx, el_ty); + auto* func = b.create( + name, params, el_ast_ty, nullptr, + ast::AttributeList{ + intrinsic, + b.Disable(ast::DisabledValidation::kFunctionHasNoBody), + }, + ast::AttributeList{}); + b.AST().AddFunction(func); + } else if (auto* arr_ty = el_ty->As()) { + // fn load_func(buffer : buf_ty, offset : u32) -> array { + // var arr : array; + // for (var i = 0u; i < array_count; i = i + 1) { + // arr[i] = el_load_func(buffer, 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(0_u)); + auto* for_init = b.Decl(i); + auto* for_cond = b.create( + ast::BinaryOp::kLessThan, b.Expr(i), b.Expr(u32(arr_ty->Count()))); + auto* for_cont = b.Assign(i, b.Add(i, 1_u)); + auto* arr_el = b.IndexAccessor(arr, i); + auto* el_offset = b.Add(b.Expr("offset"), b.Mul(i, u32(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))); - if (auto* intrinsic = IntrinsicLoadFor(ctx.dst, storage_class, el_ty)) { - auto* el_ast_ty = CreateASTTypeFor(ctx, el_ty); - auto* func = b.create( - name, params, el_ast_ty, nullptr, - ast::AttributeList{ - intrinsic, - b.Disable(ast::DisabledValidation::kFunctionHasNoBody), - }, - ast::AttributeList{}); - 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(0_u)); - auto* for_init = b.Decl(i); - auto* for_cond = b.create( - ast::BinaryOp::kLessThan, b.Expr(i), b.Expr(u32(arr_ty->Count()))); - auto* for_cont = b.Assign(i, b.Add(i, 1_u)); - auto* arr_el = b.IndexAccessor(arr, i); - auto* el_offset = b.Add(b.Expr("offset"), b.Mul(i, u32(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(buf_ty, vec_ty, var_user); - for (uint32_t i = 0; i < mat_ty->columns(); i++) { - auto* offset = b.Add("offset", u32(i * mat_ty->ColumnStride())); - values.emplace_back(b.Call(load, "buffer", offset)); - } - } else if (auto* str = el_ty->As()) { - for (auto* member : str->Members()) { - auto* offset = b.Add("offset", u32(member->Offset())); - Symbol load = LoadFunc(buf_ty, member->Type()->UnwrapRef(), var_user); - values.emplace_back(b.Call(load, "buffer", offset)); + 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(buf_ty, vec_ty, var_user); + for (uint32_t i = 0; i < mat_ty->columns(); i++) { + auto* offset = b.Add("offset", u32(i * mat_ty->ColumnStride())); + values.emplace_back(b.Call(load, "buffer", offset)); + } + } else if (auto* str = el_ty->As()) { + for (auto* member : str->Members()) { + auto* offset = b.Add("offset", u32(member->Offset())); + Symbol load = LoadFunc(buf_ty, member->Type()->UnwrapRef(), var_user); + values.emplace_back(b.Call(load, "buffer", offset)); + } } + b.Func(name, params, CreateASTTypeFor(ctx, el_ty), + { + b.Return(b.Construct(CreateASTTypeFor(ctx, el_ty), values)), + }); } - b.Func(name, params, CreateASTTypeFor(ctx, el_ty), - { - b.Return(b.Construct(CreateASTTypeFor(ctx, el_ty), values)), - }); - } - return name; - }); + return name; + }); } /// StoreFunc() returns a symbol to an intrinsic function that stores an /// 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)` + /// `fn store(buf : ptr, offset : u32, value : el_ty)` /// @param buf_ty the storage buffer type /// @param el_ty the storage buffer element type /// @param var_user the variable user @@ -523,87 +521,95 @@ struct DecomposeMemoryAccess::State { const sem::Type* el_ty, const sem::VariableUser* var_user) { auto storage_class = var_user->Variable()->StorageClass(); - return utils::GetOrCreate(store_funcs, LoadStoreKey{storage_class, buf_ty, el_ty}, [&] { - auto* buf_ast_ty = CreateASTTypeFor(ctx, buf_ty); - auto* el_ast_ty = CreateASTTypeFor(ctx, el_ty); - auto* disable_validation = - b.Disable(ast::DisabledValidation::kIgnoreConstructibleFunctionParameter); - ast::ParameterList params{ - // Note: The buffer parameter requires the StorageClass in - // order for HLSL to emit this as a ByteAddressBuffer. + auto access = var_user->Variable()->Access(); + return utils::GetOrCreate( + store_funcs, LoadStoreKey{storage_class, access, buf_ty, el_ty}, [&] { + ast::ParameterList params{ + b.Param("buffer", + b.ty.pointer(CreateASTTypeFor(ctx, buf_ty), storage_class, access), + {b.Disable(ast::DisabledValidation::kFunctionParameter)}), + b.Param("offset", b.ty.u32()), + b.Param("value", CreateASTTypeFor(ctx, el_ty)), + }; - b.create(b.Sym("buffer"), storage_class, - var_user->Variable()->Access(), buf_ast_ty, true, false, - nullptr, ast::AttributeList{disable_validation}), - b.Param("offset", b.ty.u32()), - b.Param("value", el_ast_ty), - }; + auto name = b.Sym(); - auto name = b.Sym(); + if (auto* intrinsic = IntrinsicStoreFor(ctx.dst, storage_class, el_ty)) { + auto* func = b.create( + name, params, b.ty.void_(), nullptr, + ast::AttributeList{ + intrinsic, + b.Disable(ast::DisabledValidation::kFunctionHasNoBody), + }, + ast::AttributeList{}); + b.AST().AddFunction(func); + } else { + auto body = Switch( + el_ty, // + [&](const sem::Array* arr_ty) { + // fn store_func(buffer : 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(buffer, 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(0_u)); + auto* for_init = b.Decl(i); + auto* for_cond = b.create( + ast::BinaryOp::kLessThan, b.Expr(i), b.Expr(u32(arr_ty->Count()))); + auto* for_cont = b.Assign(i, b.Add(i, 1_u)); + auto* arr_el = b.IndexAccessor(array, i); + auto* el_offset = + b.Add(b.Expr("offset"), b.Mul(i, u32(arr_ty->Stride()))); + auto* store_stmt = + b.CallStmt(b.Call(store, "buffer", el_offset, arr_el)); + auto* for_loop = + b.For(for_init, for_cond, for_cont, b.Block(store_stmt)); - if (auto* intrinsic = IntrinsicStoreFor(ctx.dst, storage_class, el_ty)) { - auto* func = b.create( - name, params, b.ty.void_(), nullptr, - ast::AttributeList{ - intrinsic, - b.Disable(ast::DisabledValidation::kFunctionHasNoBody), - }, - ast::AttributeList{}); - b.AST().AddFunction(func); - } else { - ast::StatementList body; - 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(0_u)); - auto* for_init = b.Decl(i); - auto* for_cond = b.create( - ast::BinaryOp::kLessThan, b.Expr(i), b.Expr(u32(arr_ty->Count()))); - auto* for_cont = b.Assign(i, b.Add(i, 1_u)); - auto* arr_el = b.IndexAccessor(array, i); - auto* el_offset = b.Add(b.Expr("offset"), b.Mul(i, u32(arr_ty->Stride()))); - auto* store_stmt = b.CallStmt(b.Call(store, "buffer", el_offset, arr_el)); - auto* for_loop = b.For(for_init, for_cond, for_cont, b.Block(store_stmt)); + return ast::StatementList{b.Decl(array), for_loop}; + }, + [&](const sem::Matrix* mat_ty) { + auto* vec_ty = mat_ty->ColumnType(); + Symbol store = StoreFunc(buf_ty, vec_ty, var_user); + ast::StatementList stmts; + for (uint32_t i = 0; i < mat_ty->columns(); i++) { + auto* offset = b.Add("offset", u32(i * mat_ty->ColumnStride())); + auto* element = b.IndexAccessor("value", u32(i)); + auto* call = b.Call(store, "buffer", offset, element); + stmts.emplace_back(b.CallStmt(call)); + } + return stmts; + }, + [&](const sem::Struct* str) { + ast::StatementList stmts; + for (auto* member : str->Members()) { + auto* offset = b.Add("offset", u32(member->Offset())); + auto* element = b.MemberAccessor( + "value", ctx.Clone(member->Declaration()->symbol)); + Symbol store = + StoreFunc(buf_ty, member->Type()->UnwrapRef(), var_user); + auto* call = b.Call(store, "buffer", offset, element); + stmts.emplace_back(b.CallStmt(call)); + } + return stmts; + }); - body = {b.Decl(array), for_loop}; - } else if (auto* mat_ty = el_ty->As()) { - auto* vec_ty = mat_ty->ColumnType(); - Symbol store = StoreFunc(buf_ty, vec_ty, var_user); - for (uint32_t i = 0; i < mat_ty->columns(); i++) { - auto* offset = b.Add("offset", u32(i * mat_ty->ColumnStride())); - auto* access = b.IndexAccessor("value", u32(i)); - auto* call = b.Call(store, "buffer", offset, access); - body.emplace_back(b.CallStmt(call)); - } - } else if (auto* str = el_ty->As()) { - for (auto* member : str->Members()) { - auto* offset = b.Add("offset", u32(member->Offset())); - auto* access = - b.MemberAccessor("value", ctx.Clone(member->Declaration()->symbol)); - Symbol store = StoreFunc(buf_ty, member->Type()->UnwrapRef(), var_user); - auto* call = b.Call(store, "buffer", offset, access); - body.emplace_back(b.CallStmt(call)); - } + b.Func(name, params, b.ty.void_(), body); } - b.Func(name, params, b.ty.void_(), body); - } - return name; - }); + return name; + }); } /// AtomicFunc() returns a symbol to an intrinsic function that performs an /// atomic operation from a storage buffer of type `buf_ty`. The function has /// the signature: - // `fn atomic_op(buf : buf_ty, offset : u32, ...) -> T` + // `fn atomic_op(buf : ptr, offset : u32, ...) -> T` /// @param buf_ty the storage buffer type /// @param el_ty the storage buffer element type /// @param intrinsic the atomic intrinsic @@ -614,19 +620,15 @@ struct DecomposeMemoryAccess::State { const sem::Builtin* intrinsic, const sem::VariableUser* var_user) { auto op = intrinsic->Type(); - return utils::GetOrCreate(atomic_funcs, AtomicKey{buf_ty, el_ty, op}, [&] { - auto* buf_ast_ty = CreateASTTypeFor(ctx, buf_ty); - auto* disable_validation = - b.Disable(ast::DisabledValidation::kIgnoreConstructibleFunctionParameter); + auto access = var_user->Variable()->Access(); + return utils::GetOrCreate(atomic_funcs, AtomicKey{access, buf_ty, el_ty, op}, [&] { // The first parameter to all WGSL atomics is the expression to the // atomic. This is replaced with two parameters: the buffer and offset. - ast::ParameterList params = { - // Note: The buffer parameter requires the kStorage StorageClass in - // order for HLSL to emit this as a ByteAddressBuffer. - b.create(b.Sym("buffer"), ast::StorageClass::kStorage, - var_user->Variable()->Access(), buf_ast_ty, true, false, - nullptr, ast::AttributeList{disable_validation}), + b.Param("buffer", + b.ty.pointer(CreateASTTypeFor(ctx, buf_ty), ast::StorageClass::kStorage, + access), + {b.Disable(ast::DisabledValidation::kFunctionParameter)}), b.Param("offset", b.ty.u32()), }; @@ -910,8 +912,7 @@ void DecomposeMemoryAccess::Run(CloneContext& ctx, const DataMap&, DataMap&) con if (auto* builtin = call->Target()->As()) { if (builtin->Type() == sem::BuiltinType::kArrayLength) { // arrayLength(X) - // Don't convert X into a load, this builtin actually requires the - // real pointer. + // Don't convert X into a load, this builtin actually requires the real pointer. state.TakeAccess(call_expr->args[0]); continue; } @@ -926,7 +927,7 @@ void DecomposeMemoryAccess::Run(CloneContext& ctx, const DataMap&, DataMap&) con Symbol func = state.AtomicFunc(buf_ty, el_ty, builtin, access.var->As()); - ast::ExpressionList args{ctx.Clone(buf), offset}; + ast::ExpressionList args{ctx.dst->AddressOf(ctx.Clone(buf)), offset}; for (size_t i = 1; i < call_expr->args.size(); i++) { auto* arg = call_expr->args[i]; args.emplace_back(ctx.Clone(arg)); @@ -948,26 +949,26 @@ void DecomposeMemoryAccess::Run(CloneContext& ctx, const DataMap&, DataMap&) con } BufferAccess access = access_it->second; ctx.Replace(expr, [=, &ctx, &state] { - auto* buf = access.var->Declaration(); + auto* buf = ctx.dst->AddressOf(ctx.CloneWithoutTransform(access.var->Declaration())); auto* offset = access.offset->Build(ctx); auto* buf_ty = access.var->Type()->UnwrapRef(); auto* el_ty = access.type->UnwrapRef(); Symbol func = state.LoadFunc(buf_ty, el_ty, access.var->As()); - return ctx.dst->Call(func, ctx.CloneWithoutTransform(buf), offset); + return ctx.dst->Call(func, buf, offset); }); } // And replace all storage and uniform buffer assignments with stores for (auto store : state.stores) { ctx.Replace(store.assignment, [=, &ctx, &state] { - auto* buf = store.target.var->Declaration(); + auto* buf = + ctx.dst->AddressOf(ctx.CloneWithoutTransform((store.target.var->Declaration()))); auto* offset = store.target.offset->Build(ctx); auto* buf_ty = store.target.var->Type()->UnwrapRef(); auto* el_ty = store.target.type->UnwrapRef(); auto* value = store.assignment->rhs; 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)); + auto* call = ctx.dst->Call(func, buf, offset, ctx.Clone(value)); return ctx.dst->CallStmt(call); }); } diff --git a/src/tint/transform/decompose_memory_access_test.cc b/src/tint/transform/decompose_memory_access_test.cc index 4b96bcbb42..581731e130 100644 --- a/src/tint/transform/decompose_memory_access_test.cc +++ b/src/tint/transform/decompose_memory_access_test.cc @@ -134,78 +134,78 @@ 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(@internal(disable_validation__function_parameter) buffer : ptr, 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_1(@internal(disable_validation__function_parameter) buffer : ptr, 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_2(@internal(disable_validation__function_parameter) buffer : ptr, 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_3(@internal(disable_validation__function_parameter) buffer : ptr, 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_4(@internal(disable_validation__function_parameter) buffer : ptr, 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_5(@internal(disable_validation__function_parameter) buffer : ptr, 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_6(@internal(disable_validation__function_parameter) buffer : ptr, 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_7(@internal(disable_validation__function_parameter) buffer : ptr, 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_8(@internal(disable_validation__function_parameter) buffer : ptr, 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_9(@internal(disable_validation__function_parameter) buffer : ptr, 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_10(@internal(disable_validation__function_parameter) buffer : ptr, 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_11(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32) -> vec4 -fn tint_symbol_12(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> mat2x2 { +fn tint_symbol_12(@internal(disable_validation__function_parameter) buffer : ptr, 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) -> mat2x3 { +fn tint_symbol_13(@internal(disable_validation__function_parameter) buffer : ptr, 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) -> mat2x4 { +fn tint_symbol_14(@internal(disable_validation__function_parameter) buffer : ptr, 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) -> mat3x2 { +fn tint_symbol_15(@internal(disable_validation__function_parameter) buffer : ptr, 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) -> mat3x3 { +fn tint_symbol_16(@internal(disable_validation__function_parameter) buffer : ptr, 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) -> mat3x4 { +fn tint_symbol_17(@internal(disable_validation__function_parameter) buffer : ptr, 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) -> mat4x2 { +fn tint_symbol_18(@internal(disable_validation__function_parameter) buffer : ptr, 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) -> mat4x3 { +fn tint_symbol_19(@internal(disable_validation__function_parameter) buffer : ptr, 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) -> mat4x4 { +fn tint_symbol_20(@internal(disable_validation__function_parameter) buffer : ptr, 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) -> array, 2u> { +fn tint_symbol_21(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32) -> array, 2u> { var arr : array, 2u>; for(var i_1 = 0u; (i_1 < 2u); i_1 = (i_1 + 1u)) { arr[i_1] = tint_symbol_8(buffer, (offset + (i_1 * 16u))); @@ -215,28 +215,28 @@ fn tint_symbol_21(@internal(disable_validation__ignore_constructible_function_pa @compute @workgroup_size(1) fn main() { - var a : i32 = tint_symbol(sb, 0u); - var b : u32 = tint_symbol_1(sb, 4u); - var c : f32 = tint_symbol_2(sb, 8u); - var d : vec2 = tint_symbol_3(sb, 16u); - var e : vec2 = tint_symbol_4(sb, 24u); - var f : vec2 = tint_symbol_5(sb, 32u); - var g : vec3 = tint_symbol_6(sb, 48u); - var h : vec3 = tint_symbol_7(sb, 64u); - var i : vec3 = tint_symbol_8(sb, 80u); - var j : vec4 = tint_symbol_9(sb, 96u); - var k : vec4 = tint_symbol_10(sb, 112u); - var l : vec4 = tint_symbol_11(sb, 128u); - var m : mat2x2 = tint_symbol_12(sb, 144u); - var n : mat2x3 = tint_symbol_13(sb, 160u); - var o : mat2x4 = tint_symbol_14(sb, 192u); - var p : mat3x2 = tint_symbol_15(sb, 224u); - var q : mat3x3 = tint_symbol_16(sb, 256u); - var r : mat3x4 = tint_symbol_17(sb, 304u); - var s : mat4x2 = tint_symbol_18(sb, 352u); - var t : mat4x3 = tint_symbol_19(sb, 384u); - var u : mat4x4 = tint_symbol_20(sb, 448u); - var v : array, 2> = tint_symbol_21(sb, 512u); + var a : i32 = tint_symbol(&(sb), 0u); + var b : u32 = tint_symbol_1(&(sb), 4u); + var c : f32 = tint_symbol_2(&(sb), 8u); + var d : vec2 = tint_symbol_3(&(sb), 16u); + var e : vec2 = tint_symbol_4(&(sb), 24u); + var f : vec2 = tint_symbol_5(&(sb), 32u); + var g : vec3 = tint_symbol_6(&(sb), 48u); + var h : vec3 = tint_symbol_7(&(sb), 64u); + var i : vec3 = tint_symbol_8(&(sb), 80u); + var j : vec4 = tint_symbol_9(&(sb), 96u); + var k : vec4 = tint_symbol_10(&(sb), 112u); + var l : vec4 = tint_symbol_11(&(sb), 128u); + var m : mat2x2 = tint_symbol_12(&(sb), 144u); + var n : mat2x3 = tint_symbol_13(&(sb), 160u); + var o : mat2x4 = tint_symbol_14(&(sb), 192u); + var p : mat3x2 = tint_symbol_15(&(sb), 224u); + var q : mat3x3 = tint_symbol_16(&(sb), 256u); + var r : mat3x4 = tint_symbol_17(&(sb), 304u); + var s : mat4x2 = tint_symbol_18(&(sb), 352u); + var t : mat4x3 = tint_symbol_19(&(sb), 384u); + var u : mat4x4 = tint_symbol_20(&(sb), 448u); + var v : array, 2> = tint_symbol_21(&(sb), 512u); } )"; @@ -303,78 +303,78 @@ struct SB { auto* expect = R"( @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(@internal(disable_validation__function_parameter) buffer : ptr, 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_1(@internal(disable_validation__function_parameter) buffer : ptr, 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_2(@internal(disable_validation__function_parameter) buffer : ptr, 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_3(@internal(disable_validation__function_parameter) buffer : ptr, 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_4(@internal(disable_validation__function_parameter) buffer : ptr, 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_5(@internal(disable_validation__function_parameter) buffer : ptr, 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_6(@internal(disable_validation__function_parameter) buffer : ptr, 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_7(@internal(disable_validation__function_parameter) buffer : ptr, 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_8(@internal(disable_validation__function_parameter) buffer : ptr, 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_9(@internal(disable_validation__function_parameter) buffer : ptr, 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_10(@internal(disable_validation__function_parameter) buffer : ptr, 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_11(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32) -> vec4 -fn tint_symbol_12(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> mat2x2 { +fn tint_symbol_12(@internal(disable_validation__function_parameter) buffer : ptr, 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) -> mat2x3 { +fn tint_symbol_13(@internal(disable_validation__function_parameter) buffer : ptr, 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) -> mat2x4 { +fn tint_symbol_14(@internal(disable_validation__function_parameter) buffer : ptr, 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) -> mat3x2 { +fn tint_symbol_15(@internal(disable_validation__function_parameter) buffer : ptr, 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) -> mat3x3 { +fn tint_symbol_16(@internal(disable_validation__function_parameter) buffer : ptr, 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) -> mat3x4 { +fn tint_symbol_17(@internal(disable_validation__function_parameter) buffer : ptr, 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) -> mat4x2 { +fn tint_symbol_18(@internal(disable_validation__function_parameter) buffer : ptr, 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) -> mat4x3 { +fn tint_symbol_19(@internal(disable_validation__function_parameter) buffer : ptr, 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) -> mat4x4 { +fn tint_symbol_20(@internal(disable_validation__function_parameter) buffer : ptr, 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) -> array, 2u> { +fn tint_symbol_21(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32) -> array, 2u> { var arr : array, 2u>; for(var i_1 = 0u; (i_1 < 2u); i_1 = (i_1 + 1u)) { arr[i_1] = tint_symbol_8(buffer, (offset + (i_1 * 16u))); @@ -384,28 +384,28 @@ fn tint_symbol_21(@internal(disable_validation__ignore_constructible_function_pa @compute @workgroup_size(1) fn main() { - var a : i32 = tint_symbol(sb, 0u); - var b : u32 = tint_symbol_1(sb, 4u); - var c : f32 = tint_symbol_2(sb, 8u); - var d : vec2 = tint_symbol_3(sb, 16u); - var e : vec2 = tint_symbol_4(sb, 24u); - var f : vec2 = tint_symbol_5(sb, 32u); - var g : vec3 = tint_symbol_6(sb, 48u); - var h : vec3 = tint_symbol_7(sb, 64u); - var i : vec3 = tint_symbol_8(sb, 80u); - var j : vec4 = tint_symbol_9(sb, 96u); - var k : vec4 = tint_symbol_10(sb, 112u); - var l : vec4 = tint_symbol_11(sb, 128u); - var m : mat2x2 = tint_symbol_12(sb, 144u); - var n : mat2x3 = tint_symbol_13(sb, 160u); - var o : mat2x4 = tint_symbol_14(sb, 192u); - var p : mat3x2 = tint_symbol_15(sb, 224u); - var q : mat3x3 = tint_symbol_16(sb, 256u); - var r : mat3x4 = tint_symbol_17(sb, 304u); - var s : mat4x2 = tint_symbol_18(sb, 352u); - var t : mat4x3 = tint_symbol_19(sb, 384u); - var u : mat4x4 = tint_symbol_20(sb, 448u); - var v : array, 2> = tint_symbol_21(sb, 512u); + var a : i32 = tint_symbol(&(sb), 0u); + var b : u32 = tint_symbol_1(&(sb), 4u); + var c : f32 = tint_symbol_2(&(sb), 8u); + var d : vec2 = tint_symbol_3(&(sb), 16u); + var e : vec2 = tint_symbol_4(&(sb), 24u); + var f : vec2 = tint_symbol_5(&(sb), 32u); + var g : vec3 = tint_symbol_6(&(sb), 48u); + var h : vec3 = tint_symbol_7(&(sb), 64u); + var i : vec3 = tint_symbol_8(&(sb), 80u); + var j : vec4 = tint_symbol_9(&(sb), 96u); + var k : vec4 = tint_symbol_10(&(sb), 112u); + var l : vec4 = tint_symbol_11(&(sb), 128u); + var m : mat2x2 = tint_symbol_12(&(sb), 144u); + var n : mat2x3 = tint_symbol_13(&(sb), 160u); + var o : mat2x4 = tint_symbol_14(&(sb), 192u); + var p : mat3x2 = tint_symbol_15(&(sb), 224u); + var q : mat3x3 = tint_symbol_16(&(sb), 256u); + var r : mat3x4 = tint_symbol_17(&(sb), 304u); + var s : mat4x2 = tint_symbol_18(&(sb), 352u); + var t : mat4x3 = tint_symbol_19(&(sb), 384u); + var u : mat4x4 = tint_symbol_20(&(sb), 448u); + var v : array, 2> = tint_symbol_21(&(sb), 512u); } @group(0) @binding(0) var sb : SB; @@ -526,78 +526,78 @@ struct UB { @group(0) @binding(0) var ub : UB; @internal(intrinsic_load_uniform_i32) @internal(disable_validation__function_has_no_body) -fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : UB, offset : u32) -> i32 +fn tint_symbol(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32) -> i32 @internal(intrinsic_load_uniform_u32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : UB, offset : u32) -> u32 +fn tint_symbol_1(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32) -> u32 @internal(intrinsic_load_uniform_f32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_2(@internal(disable_validation__ignore_constructible_function_parameter) buffer : UB, offset : u32) -> f32 +fn tint_symbol_2(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32) -> f32 @internal(intrinsic_load_uniform_vec2_i32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_3(@internal(disable_validation__ignore_constructible_function_parameter) buffer : UB, offset : u32) -> vec2 +fn tint_symbol_3(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32) -> vec2 @internal(intrinsic_load_uniform_vec2_u32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_4(@internal(disable_validation__ignore_constructible_function_parameter) buffer : UB, offset : u32) -> vec2 +fn tint_symbol_4(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32) -> vec2 @internal(intrinsic_load_uniform_vec2_f32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_5(@internal(disable_validation__ignore_constructible_function_parameter) buffer : UB, offset : u32) -> vec2 +fn tint_symbol_5(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32) -> vec2 @internal(intrinsic_load_uniform_vec3_i32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_6(@internal(disable_validation__ignore_constructible_function_parameter) buffer : UB, offset : u32) -> vec3 +fn tint_symbol_6(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32) -> vec3 @internal(intrinsic_load_uniform_vec3_u32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_7(@internal(disable_validation__ignore_constructible_function_parameter) buffer : UB, offset : u32) -> vec3 +fn tint_symbol_7(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32) -> vec3 @internal(intrinsic_load_uniform_vec3_f32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_8(@internal(disable_validation__ignore_constructible_function_parameter) buffer : UB, offset : u32) -> vec3 +fn tint_symbol_8(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32) -> vec3 @internal(intrinsic_load_uniform_vec4_i32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_9(@internal(disable_validation__ignore_constructible_function_parameter) buffer : UB, offset : u32) -> vec4 +fn tint_symbol_9(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32) -> vec4 @internal(intrinsic_load_uniform_vec4_u32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_10(@internal(disable_validation__ignore_constructible_function_parameter) buffer : UB, offset : u32) -> vec4 +fn tint_symbol_10(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32) -> vec4 @internal(intrinsic_load_uniform_vec4_f32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_11(@internal(disable_validation__ignore_constructible_function_parameter) buffer : UB, offset : u32) -> vec4 +fn tint_symbol_11(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32) -> vec4 -fn tint_symbol_12(@internal(disable_validation__ignore_constructible_function_parameter) buffer : UB, offset : u32) -> mat2x2 { +fn tint_symbol_12(@internal(disable_validation__function_parameter) buffer : ptr, 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 : UB, offset : u32) -> mat2x3 { +fn tint_symbol_13(@internal(disable_validation__function_parameter) buffer : ptr, 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 : UB, offset : u32) -> mat2x4 { +fn tint_symbol_14(@internal(disable_validation__function_parameter) buffer : ptr, 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 : UB, offset : u32) -> mat3x2 { +fn tint_symbol_15(@internal(disable_validation__function_parameter) buffer : ptr, 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 : UB, offset : u32) -> mat3x3 { +fn tint_symbol_16(@internal(disable_validation__function_parameter) buffer : ptr, 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 : UB, offset : u32) -> mat3x4 { +fn tint_symbol_17(@internal(disable_validation__function_parameter) buffer : ptr, 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 : UB, offset : u32) -> mat4x2 { +fn tint_symbol_18(@internal(disable_validation__function_parameter) buffer : ptr, 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 : UB, offset : u32) -> mat4x3 { +fn tint_symbol_19(@internal(disable_validation__function_parameter) buffer : ptr, 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 : UB, offset : u32) -> mat4x4 { +fn tint_symbol_20(@internal(disable_validation__function_parameter) buffer : ptr, 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 : UB, offset : u32) -> array, 2u> { +fn tint_symbol_21(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32) -> array, 2u> { var arr : array, 2u>; for(var i_1 = 0u; (i_1 < 2u); i_1 = (i_1 + 1u)) { arr[i_1] = tint_symbol_8(buffer, (offset + (i_1 * 16u))); @@ -607,28 +607,28 @@ fn tint_symbol_21(@internal(disable_validation__ignore_constructible_function_pa @compute @workgroup_size(1) fn main() { - var a : i32 = tint_symbol(ub, 0u); - var b : u32 = tint_symbol_1(ub, 4u); - var c : f32 = tint_symbol_2(ub, 8u); - var d : vec2 = tint_symbol_3(ub, 16u); - var e : vec2 = tint_symbol_4(ub, 24u); - var f : vec2 = tint_symbol_5(ub, 32u); - var g : vec3 = tint_symbol_6(ub, 48u); - var h : vec3 = tint_symbol_7(ub, 64u); - var i : vec3 = tint_symbol_8(ub, 80u); - var j : vec4 = tint_symbol_9(ub, 96u); - var k : vec4 = tint_symbol_10(ub, 112u); - var l : vec4 = tint_symbol_11(ub, 128u); - var m : mat2x2 = tint_symbol_12(ub, 144u); - var n : mat2x3 = tint_symbol_13(ub, 160u); - var o : mat2x4 = tint_symbol_14(ub, 192u); - var p : mat3x2 = tint_symbol_15(ub, 224u); - var q : mat3x3 = tint_symbol_16(ub, 256u); - var r : mat3x4 = tint_symbol_17(ub, 304u); - var s : mat4x2 = tint_symbol_18(ub, 352u); - var t : mat4x3 = tint_symbol_19(ub, 384u); - var u : mat4x4 = tint_symbol_20(ub, 448u); - var v : array, 2> = tint_symbol_21(ub, 512u); + var a : i32 = tint_symbol(&(ub), 0u); + var b : u32 = tint_symbol_1(&(ub), 4u); + var c : f32 = tint_symbol_2(&(ub), 8u); + var d : vec2 = tint_symbol_3(&(ub), 16u); + var e : vec2 = tint_symbol_4(&(ub), 24u); + var f : vec2 = tint_symbol_5(&(ub), 32u); + var g : vec3 = tint_symbol_6(&(ub), 48u); + var h : vec3 = tint_symbol_7(&(ub), 64u); + var i : vec3 = tint_symbol_8(&(ub), 80u); + var j : vec4 = tint_symbol_9(&(ub), 96u); + var k : vec4 = tint_symbol_10(&(ub), 112u); + var l : vec4 = tint_symbol_11(&(ub), 128u); + var m : mat2x2 = tint_symbol_12(&(ub), 144u); + var n : mat2x3 = tint_symbol_13(&(ub), 160u); + var o : mat2x4 = tint_symbol_14(&(ub), 192u); + var p : mat3x2 = tint_symbol_15(&(ub), 224u); + var q : mat3x3 = tint_symbol_16(&(ub), 256u); + var r : mat3x4 = tint_symbol_17(&(ub), 304u); + var s : mat4x2 = tint_symbol_18(&(ub), 352u); + var t : mat4x3 = tint_symbol_19(&(ub), 384u); + var u : mat4x4 = tint_symbol_20(&(ub), 448u); + var v : array, 2> = tint_symbol_21(&(ub), 512u); } )"; @@ -695,78 +695,78 @@ struct UB { auto* expect = R"( @internal(intrinsic_load_uniform_i32) @internal(disable_validation__function_has_no_body) -fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : UB, offset : u32) -> i32 +fn tint_symbol(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32) -> i32 @internal(intrinsic_load_uniform_u32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : UB, offset : u32) -> u32 +fn tint_symbol_1(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32) -> u32 @internal(intrinsic_load_uniform_f32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_2(@internal(disable_validation__ignore_constructible_function_parameter) buffer : UB, offset : u32) -> f32 +fn tint_symbol_2(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32) -> f32 @internal(intrinsic_load_uniform_vec2_i32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_3(@internal(disable_validation__ignore_constructible_function_parameter) buffer : UB, offset : u32) -> vec2 +fn tint_symbol_3(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32) -> vec2 @internal(intrinsic_load_uniform_vec2_u32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_4(@internal(disable_validation__ignore_constructible_function_parameter) buffer : UB, offset : u32) -> vec2 +fn tint_symbol_4(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32) -> vec2 @internal(intrinsic_load_uniform_vec2_f32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_5(@internal(disable_validation__ignore_constructible_function_parameter) buffer : UB, offset : u32) -> vec2 +fn tint_symbol_5(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32) -> vec2 @internal(intrinsic_load_uniform_vec3_i32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_6(@internal(disable_validation__ignore_constructible_function_parameter) buffer : UB, offset : u32) -> vec3 +fn tint_symbol_6(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32) -> vec3 @internal(intrinsic_load_uniform_vec3_u32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_7(@internal(disable_validation__ignore_constructible_function_parameter) buffer : UB, offset : u32) -> vec3 +fn tint_symbol_7(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32) -> vec3 @internal(intrinsic_load_uniform_vec3_f32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_8(@internal(disable_validation__ignore_constructible_function_parameter) buffer : UB, offset : u32) -> vec3 +fn tint_symbol_8(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32) -> vec3 @internal(intrinsic_load_uniform_vec4_i32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_9(@internal(disable_validation__ignore_constructible_function_parameter) buffer : UB, offset : u32) -> vec4 +fn tint_symbol_9(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32) -> vec4 @internal(intrinsic_load_uniform_vec4_u32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_10(@internal(disable_validation__ignore_constructible_function_parameter) buffer : UB, offset : u32) -> vec4 +fn tint_symbol_10(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32) -> vec4 @internal(intrinsic_load_uniform_vec4_f32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_11(@internal(disable_validation__ignore_constructible_function_parameter) buffer : UB, offset : u32) -> vec4 +fn tint_symbol_11(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32) -> vec4 -fn tint_symbol_12(@internal(disable_validation__ignore_constructible_function_parameter) buffer : UB, offset : u32) -> mat2x2 { +fn tint_symbol_12(@internal(disable_validation__function_parameter) buffer : ptr, 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 : UB, offset : u32) -> mat2x3 { +fn tint_symbol_13(@internal(disable_validation__function_parameter) buffer : ptr, 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 : UB, offset : u32) -> mat2x4 { +fn tint_symbol_14(@internal(disable_validation__function_parameter) buffer : ptr, 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 : UB, offset : u32) -> mat3x2 { +fn tint_symbol_15(@internal(disable_validation__function_parameter) buffer : ptr, 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 : UB, offset : u32) -> mat3x3 { +fn tint_symbol_16(@internal(disable_validation__function_parameter) buffer : ptr, 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 : UB, offset : u32) -> mat3x4 { +fn tint_symbol_17(@internal(disable_validation__function_parameter) buffer : ptr, 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 : UB, offset : u32) -> mat4x2 { +fn tint_symbol_18(@internal(disable_validation__function_parameter) buffer : ptr, 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 : UB, offset : u32) -> mat4x3 { +fn tint_symbol_19(@internal(disable_validation__function_parameter) buffer : ptr, 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 : UB, offset : u32) -> mat4x4 { +fn tint_symbol_20(@internal(disable_validation__function_parameter) buffer : ptr, 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 : UB, offset : u32) -> array, 2u> { +fn tint_symbol_21(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32) -> array, 2u> { var arr : array, 2u>; for(var i_1 = 0u; (i_1 < 2u); i_1 = (i_1 + 1u)) { arr[i_1] = tint_symbol_8(buffer, (offset + (i_1 * 16u))); @@ -776,28 +776,28 @@ fn tint_symbol_21(@internal(disable_validation__ignore_constructible_function_pa @compute @workgroup_size(1) fn main() { - var a : i32 = tint_symbol(ub, 0u); - var b : u32 = tint_symbol_1(ub, 4u); - var c : f32 = tint_symbol_2(ub, 8u); - var d : vec2 = tint_symbol_3(ub, 16u); - var e : vec2 = tint_symbol_4(ub, 24u); - var f : vec2 = tint_symbol_5(ub, 32u); - var g : vec3 = tint_symbol_6(ub, 48u); - var h : vec3 = tint_symbol_7(ub, 64u); - var i : vec3 = tint_symbol_8(ub, 80u); - var j : vec4 = tint_symbol_9(ub, 96u); - var k : vec4 = tint_symbol_10(ub, 112u); - var l : vec4 = tint_symbol_11(ub, 128u); - var m : mat2x2 = tint_symbol_12(ub, 144u); - var n : mat2x3 = tint_symbol_13(ub, 160u); - var o : mat2x4 = tint_symbol_14(ub, 192u); - var p : mat3x2 = tint_symbol_15(ub, 224u); - var q : mat3x3 = tint_symbol_16(ub, 256u); - var r : mat3x4 = tint_symbol_17(ub, 304u); - var s : mat4x2 = tint_symbol_18(ub, 352u); - var t : mat4x3 = tint_symbol_19(ub, 384u); - var u : mat4x4 = tint_symbol_20(ub, 448u); - var v : array, 2> = tint_symbol_21(ub, 512u); + var a : i32 = tint_symbol(&(ub), 0u); + var b : u32 = tint_symbol_1(&(ub), 4u); + var c : f32 = tint_symbol_2(&(ub), 8u); + var d : vec2 = tint_symbol_3(&(ub), 16u); + var e : vec2 = tint_symbol_4(&(ub), 24u); + var f : vec2 = tint_symbol_5(&(ub), 32u); + var g : vec3 = tint_symbol_6(&(ub), 48u); + var h : vec3 = tint_symbol_7(&(ub), 64u); + var i : vec3 = tint_symbol_8(&(ub), 80u); + var j : vec4 = tint_symbol_9(&(ub), 96u); + var k : vec4 = tint_symbol_10(&(ub), 112u); + var l : vec4 = tint_symbol_11(&(ub), 128u); + var m : mat2x2 = tint_symbol_12(&(ub), 144u); + var n : mat2x3 = tint_symbol_13(&(ub), 160u); + var o : mat2x4 = tint_symbol_14(&(ub), 192u); + var p : mat3x2 = tint_symbol_15(&(ub), 224u); + var q : mat3x3 = tint_symbol_16(&(ub), 256u); + var r : mat3x4 = tint_symbol_17(&(ub), 304u); + var s : mat4x2 = tint_symbol_18(&(ub), 352u); + var t : mat4x3 = tint_symbol_19(&(ub), 384u); + var u : mat4x4 = tint_symbol_20(&(ub), 448u); + var v : array, 2> = tint_symbol_21(&(ub), 512u); } @group(0) @binding(0) var ub : UB; @@ -918,96 +918,96 @@ 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(@internal(disable_validation__function_parameter) buffer : ptr, 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_1(@internal(disable_validation__function_parameter) buffer : ptr, 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_2(@internal(disable_validation__function_parameter) buffer : ptr, 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_3(@internal(disable_validation__function_parameter) buffer : ptr, 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_4(@internal(disable_validation__function_parameter) buffer : ptr, 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_5(@internal(disable_validation__function_parameter) buffer : ptr, 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_6(@internal(disable_validation__function_parameter) buffer : ptr, 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_7(@internal(disable_validation__function_parameter) buffer : ptr, 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_8(@internal(disable_validation__function_parameter) buffer : ptr, 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_9(@internal(disable_validation__function_parameter) buffer : ptr, 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_10(@internal(disable_validation__function_parameter) buffer : ptr, 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_11(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, value : vec4) -fn tint_symbol_12(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : mat2x2) { +fn tint_symbol_12(@internal(disable_validation__function_parameter) buffer : ptr, 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 : mat2x3) { +fn tint_symbol_13(@internal(disable_validation__function_parameter) buffer : ptr, 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 : mat2x4) { +fn tint_symbol_14(@internal(disable_validation__function_parameter) buffer : ptr, 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 : mat3x2) { +fn tint_symbol_15(@internal(disable_validation__function_parameter) buffer : ptr, 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 : mat3x3) { +fn tint_symbol_16(@internal(disable_validation__function_parameter) buffer : ptr, 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 : mat3x4) { +fn tint_symbol_17(@internal(disable_validation__function_parameter) buffer : ptr, 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 : mat4x2) { +fn tint_symbol_18(@internal(disable_validation__function_parameter) buffer : ptr, 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 : mat4x3) { +fn tint_symbol_19(@internal(disable_validation__function_parameter) buffer : ptr, 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 : mat4x4) { +fn tint_symbol_20(@internal(disable_validation__function_parameter) buffer : ptr, 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 : array, 2u>) { +fn tint_symbol_21(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, value : array, 2u>) { 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]); @@ -1016,28 +1016,28 @@ fn tint_symbol_21(@internal(disable_validation__ignore_constructible_function_pa @compute @workgroup_size(1) fn main() { - tint_symbol(sb, 0u, i32()); - tint_symbol_1(sb, 4u, u32()); - tint_symbol_2(sb, 8u, f32()); - tint_symbol_3(sb, 16u, vec2()); - tint_symbol_4(sb, 24u, vec2()); - tint_symbol_5(sb, 32u, vec2()); - tint_symbol_6(sb, 48u, vec3()); - tint_symbol_7(sb, 64u, vec3()); - tint_symbol_8(sb, 80u, vec3()); - tint_symbol_9(sb, 96u, vec4()); - tint_symbol_10(sb, 112u, vec4()); - tint_symbol_11(sb, 128u, vec4()); - tint_symbol_12(sb, 144u, mat2x2()); - tint_symbol_13(sb, 160u, mat2x3()); - tint_symbol_14(sb, 192u, mat2x4()); - tint_symbol_15(sb, 224u, mat3x2()); - tint_symbol_16(sb, 256u, mat3x3()); - tint_symbol_17(sb, 304u, mat3x4()); - tint_symbol_18(sb, 352u, mat4x2()); - tint_symbol_19(sb, 384u, mat4x3()); - tint_symbol_20(sb, 448u, mat4x4()); - tint_symbol_21(sb, 512u, array, 2>()); + tint_symbol(&(sb), 0u, i32()); + tint_symbol_1(&(sb), 4u, u32()); + tint_symbol_2(&(sb), 8u, f32()); + tint_symbol_3(&(sb), 16u, vec2()); + tint_symbol_4(&(sb), 24u, vec2()); + tint_symbol_5(&(sb), 32u, vec2()); + tint_symbol_6(&(sb), 48u, vec3()); + tint_symbol_7(&(sb), 64u, vec3()); + tint_symbol_8(&(sb), 80u, vec3()); + tint_symbol_9(&(sb), 96u, vec4()); + tint_symbol_10(&(sb), 112u, vec4()); + tint_symbol_11(&(sb), 128u, vec4()); + tint_symbol_12(&(sb), 144u, mat2x2()); + tint_symbol_13(&(sb), 160u, mat2x3()); + tint_symbol_14(&(sb), 192u, mat2x4()); + tint_symbol_15(&(sb), 224u, mat3x2()); + tint_symbol_16(&(sb), 256u, mat3x3()); + tint_symbol_17(&(sb), 304u, mat3x4()); + tint_symbol_18(&(sb), 352u, mat4x2()); + tint_symbol_19(&(sb), 384u, mat4x3()); + tint_symbol_20(&(sb), 448u, mat4x4()); + tint_symbol_21(&(sb), 512u, array, 2>()); } )"; @@ -1104,96 +1104,96 @@ struct SB { auto* expect = R"( @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(@internal(disable_validation__function_parameter) buffer : ptr, 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_1(@internal(disable_validation__function_parameter) buffer : ptr, 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_2(@internal(disable_validation__function_parameter) buffer : ptr, 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_3(@internal(disable_validation__function_parameter) buffer : ptr, 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_4(@internal(disable_validation__function_parameter) buffer : ptr, 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_5(@internal(disable_validation__function_parameter) buffer : ptr, 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_6(@internal(disable_validation__function_parameter) buffer : ptr, 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_7(@internal(disable_validation__function_parameter) buffer : ptr, 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_8(@internal(disable_validation__function_parameter) buffer : ptr, 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_9(@internal(disable_validation__function_parameter) buffer : ptr, 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_10(@internal(disable_validation__function_parameter) buffer : ptr, 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_11(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, value : vec4) -fn tint_symbol_12(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : mat2x2) { +fn tint_symbol_12(@internal(disable_validation__function_parameter) buffer : ptr, 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 : mat2x3) { +fn tint_symbol_13(@internal(disable_validation__function_parameter) buffer : ptr, 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 : mat2x4) { +fn tint_symbol_14(@internal(disable_validation__function_parameter) buffer : ptr, 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 : mat3x2) { +fn tint_symbol_15(@internal(disable_validation__function_parameter) buffer : ptr, 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 : mat3x3) { +fn tint_symbol_16(@internal(disable_validation__function_parameter) buffer : ptr, 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 : mat3x4) { +fn tint_symbol_17(@internal(disable_validation__function_parameter) buffer : ptr, 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 : mat4x2) { +fn tint_symbol_18(@internal(disable_validation__function_parameter) buffer : ptr, 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 : mat4x3) { +fn tint_symbol_19(@internal(disable_validation__function_parameter) buffer : ptr, 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 : mat4x4) { +fn tint_symbol_20(@internal(disable_validation__function_parameter) buffer : ptr, 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 : array, 2u>) { +fn tint_symbol_21(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, value : array, 2u>) { 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]); @@ -1202,28 +1202,28 @@ fn tint_symbol_21(@internal(disable_validation__ignore_constructible_function_pa @compute @workgroup_size(1) fn main() { - tint_symbol(sb, 0u, i32()); - tint_symbol_1(sb, 4u, u32()); - tint_symbol_2(sb, 8u, f32()); - tint_symbol_3(sb, 16u, vec2()); - tint_symbol_4(sb, 24u, vec2()); - tint_symbol_5(sb, 32u, vec2()); - tint_symbol_6(sb, 48u, vec3()); - tint_symbol_7(sb, 64u, vec3()); - tint_symbol_8(sb, 80u, vec3()); - tint_symbol_9(sb, 96u, vec4()); - tint_symbol_10(sb, 112u, vec4()); - tint_symbol_11(sb, 128u, vec4()); - tint_symbol_12(sb, 144u, mat2x2()); - tint_symbol_13(sb, 160u, mat2x3()); - tint_symbol_14(sb, 192u, mat2x4()); - tint_symbol_15(sb, 224u, mat3x2()); - tint_symbol_16(sb, 256u, mat3x3()); - tint_symbol_17(sb, 304u, mat3x4()); - tint_symbol_18(sb, 352u, mat4x2()); - tint_symbol_19(sb, 384u, mat4x3()); - tint_symbol_20(sb, 448u, mat4x4()); - tint_symbol_21(sb, 512u, array, 2>()); + tint_symbol(&(sb), 0u, i32()); + tint_symbol_1(&(sb), 4u, u32()); + tint_symbol_2(&(sb), 8u, f32()); + tint_symbol_3(&(sb), 16u, vec2()); + tint_symbol_4(&(sb), 24u, vec2()); + tint_symbol_5(&(sb), 32u, vec2()); + tint_symbol_6(&(sb), 48u, vec3()); + tint_symbol_7(&(sb), 64u, vec3()); + tint_symbol_8(&(sb), 80u, vec3()); + tint_symbol_9(&(sb), 96u, vec4()); + tint_symbol_10(&(sb), 112u, vec4()); + tint_symbol_11(&(sb), 128u, vec4()); + tint_symbol_12(&(sb), 144u, mat2x2()); + tint_symbol_13(&(sb), 160u, mat2x3()); + tint_symbol_14(&(sb), 192u, mat2x4()); + tint_symbol_15(&(sb), 224u, mat3x2()); + tint_symbol_16(&(sb), 256u, mat3x3()); + tint_symbol_17(&(sb), 304u, mat3x4()); + tint_symbol_18(&(sb), 352u, mat4x2()); + tint_symbol_19(&(sb), 384u, mat4x3()); + tint_symbol_20(&(sb), 448u, mat4x4()); + tint_symbol_21(&(sb), 512u, array, 2>()); } @group(0) @binding(0) var sb : SB; @@ -1323,78 +1323,78 @@ struct SB { @group(0) @binding(0) var sb : SB; @internal(intrinsic_load_storage_i32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> i32 +fn tint_symbol_1(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32) -> i32 @internal(intrinsic_load_storage_u32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_2(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> u32 +fn tint_symbol_2(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32) -> u32 @internal(intrinsic_load_storage_f32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_3(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> f32 +fn tint_symbol_3(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32) -> f32 @internal(intrinsic_load_storage_vec2_i32) @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_4(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32) -> vec2 @internal(intrinsic_load_storage_vec2_u32) @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_5(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32) -> vec2 @internal(intrinsic_load_storage_vec2_f32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_6(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> vec2 +fn tint_symbol_6(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32) -> vec2 @internal(intrinsic_load_storage_vec3_i32) @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_7(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32) -> vec3 @internal(intrinsic_load_storage_vec3_u32) @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_8(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32) -> vec3 @internal(intrinsic_load_storage_vec3_f32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_9(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> vec3 +fn tint_symbol_9(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32) -> vec3 @internal(intrinsic_load_storage_vec4_i32) @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_10(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32) -> vec4 @internal(intrinsic_load_storage_vec4_u32) @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_11(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32) -> vec4 @internal(intrinsic_load_storage_vec4_f32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_12(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> vec4 +fn tint_symbol_12(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32) -> vec4 -fn tint_symbol_13(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> mat2x2 { +fn tint_symbol_13(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32) -> mat2x2 { return mat2x2(tint_symbol_6(buffer, (offset + 0u)), tint_symbol_6(buffer, (offset + 8u))); } -fn tint_symbol_14(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> mat2x3 { +fn tint_symbol_14(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32) -> mat2x3 { return mat2x3(tint_symbol_9(buffer, (offset + 0u)), tint_symbol_9(buffer, (offset + 16u))); } -fn tint_symbol_15(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> mat2x4 { +fn tint_symbol_15(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32) -> mat2x4 { return mat2x4(tint_symbol_12(buffer, (offset + 0u)), tint_symbol_12(buffer, (offset + 16u))); } -fn tint_symbol_16(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> mat3x2 { +fn tint_symbol_16(@internal(disable_validation__function_parameter) buffer : ptr, 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_17(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> mat3x3 { +fn tint_symbol_17(@internal(disable_validation__function_parameter) buffer : ptr, 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_18(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> mat3x4 { +fn tint_symbol_18(@internal(disable_validation__function_parameter) buffer : ptr, 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_19(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> mat4x2 { +fn tint_symbol_19(@internal(disable_validation__function_parameter) buffer : ptr, 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_20(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> mat4x3 { +fn tint_symbol_20(@internal(disable_validation__function_parameter) buffer : ptr, 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_21(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> mat4x4 { +fn tint_symbol_21(@internal(disable_validation__function_parameter) buffer : ptr, 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_22(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> array, 2u> { +fn tint_symbol_22(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32) -> array, 2u> { var arr : array, 2u>; for(var i_1 = 0u; (i_1 < 2u); i_1 = (i_1 + 1u)) { arr[i_1] = tint_symbol_9(buffer, (offset + (i_1 * 16u))); @@ -1402,13 +1402,13 @@ fn tint_symbol_22(@internal(disable_validation__ignore_constructible_function_pa return arr; } -fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> SB { +fn tint_symbol(@internal(disable_validation__function_parameter) buffer : ptr, 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))); } @compute @workgroup_size(1) fn main() { - var x : SB = tint_symbol(sb, 0u); + var x : SB = tint_symbol(&(sb), 0u); } )"; @@ -1454,78 +1454,78 @@ struct SB { auto* expect = R"( @internal(intrinsic_load_storage_i32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> i32 +fn tint_symbol_1(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32) -> i32 @internal(intrinsic_load_storage_u32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_2(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> u32 +fn tint_symbol_2(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32) -> u32 @internal(intrinsic_load_storage_f32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_3(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> f32 +fn tint_symbol_3(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32) -> f32 @internal(intrinsic_load_storage_vec2_i32) @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_4(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32) -> vec2 @internal(intrinsic_load_storage_vec2_u32) @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_5(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32) -> vec2 @internal(intrinsic_load_storage_vec2_f32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_6(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> vec2 +fn tint_symbol_6(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32) -> vec2 @internal(intrinsic_load_storage_vec3_i32) @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_7(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32) -> vec3 @internal(intrinsic_load_storage_vec3_u32) @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_8(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32) -> vec3 @internal(intrinsic_load_storage_vec3_f32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_9(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> vec3 +fn tint_symbol_9(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32) -> vec3 @internal(intrinsic_load_storage_vec4_i32) @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_10(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32) -> vec4 @internal(intrinsic_load_storage_vec4_u32) @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_11(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32) -> vec4 @internal(intrinsic_load_storage_vec4_f32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_12(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> vec4 +fn tint_symbol_12(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32) -> vec4 -fn tint_symbol_13(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> mat2x2 { +fn tint_symbol_13(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32) -> mat2x2 { return mat2x2(tint_symbol_6(buffer, (offset + 0u)), tint_symbol_6(buffer, (offset + 8u))); } -fn tint_symbol_14(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> mat2x3 { +fn tint_symbol_14(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32) -> mat2x3 { return mat2x3(tint_symbol_9(buffer, (offset + 0u)), tint_symbol_9(buffer, (offset + 16u))); } -fn tint_symbol_15(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> mat2x4 { +fn tint_symbol_15(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32) -> mat2x4 { return mat2x4(tint_symbol_12(buffer, (offset + 0u)), tint_symbol_12(buffer, (offset + 16u))); } -fn tint_symbol_16(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> mat3x2 { +fn tint_symbol_16(@internal(disable_validation__function_parameter) buffer : ptr, 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_17(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> mat3x3 { +fn tint_symbol_17(@internal(disable_validation__function_parameter) buffer : ptr, 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_18(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> mat3x4 { +fn tint_symbol_18(@internal(disable_validation__function_parameter) buffer : ptr, 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_19(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> mat4x2 { +fn tint_symbol_19(@internal(disable_validation__function_parameter) buffer : ptr, 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_20(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> mat4x3 { +fn tint_symbol_20(@internal(disable_validation__function_parameter) buffer : ptr, 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_21(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> mat4x4 { +fn tint_symbol_21(@internal(disable_validation__function_parameter) buffer : ptr, 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_22(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> array, 2u> { +fn tint_symbol_22(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32) -> array, 2u> { var arr : array, 2u>; for(var i_1 = 0u; (i_1 < 2u); i_1 = (i_1 + 1u)) { arr[i_1] = tint_symbol_9(buffer, (offset + (i_1 * 16u))); @@ -1533,13 +1533,13 @@ fn tint_symbol_22(@internal(disable_validation__ignore_constructible_function_pa return arr; } -fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> SB { +fn tint_symbol(@internal(disable_validation__function_parameter) buffer : ptr, 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))); } @compute @workgroup_size(1) fn main() { - var x : SB = tint_symbol(sb, 0u); + var x : SB = tint_symbol(&(sb), 0u); } @group(0) @binding(0) var sb : SB; @@ -1639,103 +1639,103 @@ struct SB { @group(0) @binding(0) var sb : SB; @internal(intrinsic_store_storage_i32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : i32) +fn tint_symbol_1(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, value : i32) @internal(intrinsic_store_storage_u32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_2(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : u32) +fn tint_symbol_2(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, value : u32) @internal(intrinsic_store_storage_f32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_3(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : f32) +fn tint_symbol_3(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, value : f32) @internal(intrinsic_store_storage_vec2_i32) @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_4(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, value : vec2) @internal(intrinsic_store_storage_vec2_u32) @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_5(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, value : vec2) @internal(intrinsic_store_storage_vec2_f32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_6(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : vec2) +fn tint_symbol_6(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, value : vec2) @internal(intrinsic_store_storage_vec3_i32) @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_7(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, value : vec3) @internal(intrinsic_store_storage_vec3_u32) @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_8(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, value : vec3) @internal(intrinsic_store_storage_vec3_f32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_9(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : vec3) +fn tint_symbol_9(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, value : vec3) @internal(intrinsic_store_storage_vec4_i32) @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_10(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, value : vec4) @internal(intrinsic_store_storage_vec4_u32) @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_11(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, value : vec4) @internal(intrinsic_store_storage_vec4_f32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_12(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : vec4) +fn tint_symbol_12(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, value : vec4) -fn tint_symbol_13(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : mat2x2) { +fn tint_symbol_13(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, value : mat2x2) { tint_symbol_6(buffer, (offset + 0u), value[0u]); tint_symbol_6(buffer, (offset + 8u), value[1u]); } -fn tint_symbol_14(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : mat2x3) { +fn tint_symbol_14(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, value : mat2x3) { tint_symbol_9(buffer, (offset + 0u), value[0u]); tint_symbol_9(buffer, (offset + 16u), value[1u]); } -fn tint_symbol_15(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : mat2x4) { +fn tint_symbol_15(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, value : mat2x4) { tint_symbol_12(buffer, (offset + 0u), value[0u]); tint_symbol_12(buffer, (offset + 16u), value[1u]); } -fn tint_symbol_16(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : mat3x2) { +fn tint_symbol_16(@internal(disable_validation__function_parameter) buffer : ptr, 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_17(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : mat3x3) { +fn tint_symbol_17(@internal(disable_validation__function_parameter) buffer : ptr, 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_18(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : mat3x4) { +fn tint_symbol_18(@internal(disable_validation__function_parameter) buffer : ptr, 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_19(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : mat4x2) { +fn tint_symbol_19(@internal(disable_validation__function_parameter) buffer : ptr, 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_20(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : mat4x3) { +fn tint_symbol_20(@internal(disable_validation__function_parameter) buffer : ptr, 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_21(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : mat4x4) { +fn tint_symbol_21(@internal(disable_validation__function_parameter) buffer : ptr, 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_22(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : array, 2u>) { +fn tint_symbol_22(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, value : array, 2u>) { 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(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : SB) { +fn tint_symbol(@internal(disable_validation__function_parameter) buffer : ptr, 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); @@ -1762,7 +1762,7 @@ fn tint_symbol(@internal(disable_validation__ignore_constructible_function_param @compute @workgroup_size(1) fn main() { - tint_symbol(sb, 0u, SB()); + tint_symbol(&(sb), 0u, SB()); } )"; @@ -1808,103 +1808,103 @@ struct SB { auto* expect = R"( @internal(intrinsic_store_storage_i32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : i32) +fn tint_symbol_1(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, value : i32) @internal(intrinsic_store_storage_u32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_2(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : u32) +fn tint_symbol_2(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, value : u32) @internal(intrinsic_store_storage_f32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_3(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : f32) +fn tint_symbol_3(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, value : f32) @internal(intrinsic_store_storage_vec2_i32) @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_4(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, value : vec2) @internal(intrinsic_store_storage_vec2_u32) @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_5(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, value : vec2) @internal(intrinsic_store_storage_vec2_f32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_6(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : vec2) +fn tint_symbol_6(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, value : vec2) @internal(intrinsic_store_storage_vec3_i32) @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_7(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, value : vec3) @internal(intrinsic_store_storage_vec3_u32) @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_8(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, value : vec3) @internal(intrinsic_store_storage_vec3_f32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_9(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : vec3) +fn tint_symbol_9(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, value : vec3) @internal(intrinsic_store_storage_vec4_i32) @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_10(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, value : vec4) @internal(intrinsic_store_storage_vec4_u32) @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_11(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, value : vec4) @internal(intrinsic_store_storage_vec4_f32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_12(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : vec4) +fn tint_symbol_12(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, value : vec4) -fn tint_symbol_13(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : mat2x2) { +fn tint_symbol_13(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, value : mat2x2) { tint_symbol_6(buffer, (offset + 0u), value[0u]); tint_symbol_6(buffer, (offset + 8u), value[1u]); } -fn tint_symbol_14(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : mat2x3) { +fn tint_symbol_14(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, value : mat2x3) { tint_symbol_9(buffer, (offset + 0u), value[0u]); tint_symbol_9(buffer, (offset + 16u), value[1u]); } -fn tint_symbol_15(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : mat2x4) { +fn tint_symbol_15(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, value : mat2x4) { tint_symbol_12(buffer, (offset + 0u), value[0u]); tint_symbol_12(buffer, (offset + 16u), value[1u]); } -fn tint_symbol_16(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : mat3x2) { +fn tint_symbol_16(@internal(disable_validation__function_parameter) buffer : ptr, 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_17(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : mat3x3) { +fn tint_symbol_17(@internal(disable_validation__function_parameter) buffer : ptr, 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_18(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : mat3x4) { +fn tint_symbol_18(@internal(disable_validation__function_parameter) buffer : ptr, 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_19(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : mat4x2) { +fn tint_symbol_19(@internal(disable_validation__function_parameter) buffer : ptr, 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_20(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : mat4x3) { +fn tint_symbol_20(@internal(disable_validation__function_parameter) buffer : ptr, 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_21(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : mat4x4) { +fn tint_symbol_21(@internal(disable_validation__function_parameter) buffer : ptr, 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_22(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : array, 2u>) { +fn tint_symbol_22(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, value : array, 2u>) { 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(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : SB) { +fn tint_symbol(@internal(disable_validation__function_parameter) buffer : ptr, 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); @@ -1931,7 +1931,7 @@ fn tint_symbol(@internal(disable_validation__ignore_constructible_function_param @compute @workgroup_size(1) fn main() { - tint_symbol(sb, 0u, SB()); + tint_symbol(&(sb), 0u, SB()); } @group(0) @binding(0) var sb : SB; @@ -2028,11 +2028,11 @@ struct SB { @group(0) @binding(0) var sb : SB; @internal(intrinsic_load_storage_f32) @internal(disable_validation__function_has_no_body) -fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> f32 +fn tint_symbol(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32) -> f32 @compute @workgroup_size(1) fn main() { - var x : f32 = tint_symbol(sb, 712u); + var x : f32 = tint_symbol(&(sb), 712u); } )"; @@ -2078,11 +2078,11 @@ struct S1 { auto* expect = R"( @internal(intrinsic_load_storage_f32) @internal(disable_validation__function_has_no_body) -fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> f32 +fn tint_symbol(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32) -> f32 @compute @workgroup_size(1) fn main() { - var x : f32 = tint_symbol(sb, 712u); + var x : f32 = tint_symbol(&(sb), 712u); } @group(0) @binding(0) var sb : SB; @@ -2164,14 +2164,14 @@ struct SB { @group(0) @binding(0) var sb : SB; @internal(intrinsic_load_storage_f32) @internal(disable_validation__function_has_no_body) -fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> f32 +fn tint_symbol(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32) -> f32 @compute @workgroup_size(1) fn main() { var i : i32 = 4; var j : u32 = 1u; var k : i32 = 2; - var x : f32 = tint_symbol(sb, (((((128u + (128u * u32(i))) + 16u) + (32u * j)) + 16u) + (4u * u32(k)))); + var x : f32 = tint_symbol(&(sb), (((((128u + (128u * u32(i))) + 16u) + (32u * j)) + 16u) + (4u * u32(k)))); } )"; @@ -2213,14 +2213,14 @@ struct S1 { auto* expect = R"( @internal(intrinsic_load_storage_f32) @internal(disable_validation__function_has_no_body) -fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> f32 +fn tint_symbol(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32) -> f32 @compute @workgroup_size(1) fn main() { var i : i32 = 4; var j : u32 = 1u; var k : i32 = 2; - var x : f32 = tint_symbol(sb, (((((128u + (128u * u32(i))) + 16u) + (32u * j)) + 16u) + (4u * u32(k)))); + var x : f32 = tint_symbol(&(sb), (((((128u + (128u * u32(i))) + 16u) + (32u * j)) + 16u) + (4u * u32(k)))); } @group(0) @binding(0) var sb : SB; @@ -2318,14 +2318,14 @@ struct SB { @group(0) @binding(0) var sb : SB; @internal(intrinsic_load_storage_f32) @internal(disable_validation__function_has_no_body) -fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> f32 +fn tint_symbol(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32) -> f32 @compute @workgroup_size(1) fn main() { var i : i32 = 4; var j : u32 = 1u; var k : i32 = 2; - var x : f32 = tint_symbol(sb, (((((128u + (128u * u32(i))) + 16u) + (32u * j)) + 16u) + (4u * u32(k)))); + var x : f32 = tint_symbol(&(sb), (((((128u + (128u * u32(i))) + 16u) + (32u * j)) + 16u) + (4u * u32(k)))); } )"; @@ -2375,14 +2375,14 @@ struct S1 { auto* expect = R"( @internal(intrinsic_load_storage_f32) @internal(disable_validation__function_has_no_body) -fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> f32 +fn tint_symbol(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32) -> f32 @compute @workgroup_size(1) fn main() { var i : i32 = 4; var j : u32 = 1u; var k : i32 = 2; - var x : f32 = tint_symbol(sb, (((((128u + (128u * u32(i))) + 16u) + (32u * j)) + 16u) + (4u * u32(k)))); + var x : f32 = tint_symbol(&(sb), (((((128u + (128u * u32(i))) + 16u) + (32u * j)) + 16u) + (4u * u32(k)))); } @group(0) @binding(0) var sb : SB; @@ -2467,34 +2467,34 @@ struct SB { @group(0) @binding(0) var sb : SB; @internal(intrinsic_atomic_store_storage_i32) @internal(disable_validation__function_has_no_body) -fn tint_atomicStore(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) +fn tint_atomicStore(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, param_1 : i32) @internal(intrinsic_atomic_load_storage_i32) @internal(disable_validation__function_has_no_body) -fn tint_atomicLoad(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> i32 +fn tint_atomicLoad(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32) -> i32 @internal(intrinsic_atomic_add_storage_i32) @internal(disable_validation__function_has_no_body) -fn tint_atomicAdd(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32 +fn tint_atomicAdd(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, param_1 : i32) -> i32 @internal(intrinsic_atomic_sub_storage_i32) @internal(disable_validation__function_has_no_body) -fn tint_atomicSub(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32 +fn tint_atomicSub(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, param_1 : i32) -> i32 @internal(intrinsic_atomic_max_storage_i32) @internal(disable_validation__function_has_no_body) -fn tint_atomicMax(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32 +fn tint_atomicMax(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, param_1 : i32) -> i32 @internal(intrinsic_atomic_min_storage_i32) @internal(disable_validation__function_has_no_body) -fn tint_atomicMin(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32 +fn tint_atomicMin(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, param_1 : i32) -> i32 @internal(intrinsic_atomic_and_storage_i32) @internal(disable_validation__function_has_no_body) -fn tint_atomicAnd(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32 +fn tint_atomicAnd(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, param_1 : i32) -> i32 @internal(intrinsic_atomic_or_storage_i32) @internal(disable_validation__function_has_no_body) -fn tint_atomicOr(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32 +fn tint_atomicOr(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, param_1 : i32) -> i32 @internal(intrinsic_atomic_xor_storage_i32) @internal(disable_validation__function_has_no_body) -fn tint_atomicXor(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32 +fn tint_atomicXor(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, param_1 : i32) -> i32 @internal(intrinsic_atomic_exchange_storage_i32) @internal(disable_validation__function_has_no_body) -fn tint_atomicExchange(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32 +fn tint_atomicExchange(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, param_1 : i32) -> i32 struct atomic_compare_exchange_weak_ret_type { old_value : i32, @@ -2502,37 +2502,37 @@ struct atomic_compare_exchange_weak_ret_type { } @internal(intrinsic_atomic_compare_exchange_weak_storage_i32) @internal(disable_validation__function_has_no_body) -fn tint_atomicCompareExchangeWeak(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32, param_2 : i32) -> atomic_compare_exchange_weak_ret_type +fn tint_atomicCompareExchangeWeak(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, param_1 : i32, param_2 : i32) -> atomic_compare_exchange_weak_ret_type @internal(intrinsic_atomic_store_storage_u32) @internal(disable_validation__function_has_no_body) -fn tint_atomicStore_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) +fn tint_atomicStore_1(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, param_1 : u32) @internal(intrinsic_atomic_load_storage_u32) @internal(disable_validation__function_has_no_body) -fn tint_atomicLoad_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> u32 +fn tint_atomicLoad_1(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32) -> u32 @internal(intrinsic_atomic_add_storage_u32) @internal(disable_validation__function_has_no_body) -fn tint_atomicAdd_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32 +fn tint_atomicAdd_1(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, param_1 : u32) -> u32 @internal(intrinsic_atomic_sub_storage_u32) @internal(disable_validation__function_has_no_body) -fn tint_atomicSub_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32 +fn tint_atomicSub_1(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, param_1 : u32) -> u32 @internal(intrinsic_atomic_max_storage_u32) @internal(disable_validation__function_has_no_body) -fn tint_atomicMax_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32 +fn tint_atomicMax_1(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, param_1 : u32) -> u32 @internal(intrinsic_atomic_min_storage_u32) @internal(disable_validation__function_has_no_body) -fn tint_atomicMin_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32 +fn tint_atomicMin_1(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, param_1 : u32) -> u32 @internal(intrinsic_atomic_and_storage_u32) @internal(disable_validation__function_has_no_body) -fn tint_atomicAnd_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32 +fn tint_atomicAnd_1(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, param_1 : u32) -> u32 @internal(intrinsic_atomic_or_storage_u32) @internal(disable_validation__function_has_no_body) -fn tint_atomicOr_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32 +fn tint_atomicOr_1(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, param_1 : u32) -> u32 @internal(intrinsic_atomic_xor_storage_u32) @internal(disable_validation__function_has_no_body) -fn tint_atomicXor_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32 +fn tint_atomicXor_1(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, param_1 : u32) -> u32 @internal(intrinsic_atomic_exchange_storage_u32) @internal(disable_validation__function_has_no_body) -fn tint_atomicExchange_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32 +fn tint_atomicExchange_1(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, param_1 : u32) -> u32 struct atomic_compare_exchange_weak_ret_type_1 { old_value : u32, @@ -2540,32 +2540,32 @@ struct atomic_compare_exchange_weak_ret_type_1 { } @internal(intrinsic_atomic_compare_exchange_weak_storage_u32) @internal(disable_validation__function_has_no_body) -fn tint_atomicCompareExchangeWeak_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32, param_2 : u32) -> atomic_compare_exchange_weak_ret_type_1 +fn tint_atomicCompareExchangeWeak_1(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, param_1 : u32, param_2 : u32) -> atomic_compare_exchange_weak_ret_type_1 @compute @workgroup_size(1) fn main() { - tint_atomicStore(sb, 16u, 123); - tint_atomicLoad(sb, 16u); - tint_atomicAdd(sb, 16u, 123); - tint_atomicSub(sb, 16u, 123); - tint_atomicMax(sb, 16u, 123); - tint_atomicMin(sb, 16u, 123); - tint_atomicAnd(sb, 16u, 123); - tint_atomicOr(sb, 16u, 123); - tint_atomicXor(sb, 16u, 123); - tint_atomicExchange(sb, 16u, 123); - tint_atomicCompareExchangeWeak(sb, 16u, 123, 345); - tint_atomicStore_1(sb, 20u, 123u); - tint_atomicLoad_1(sb, 20u); - tint_atomicAdd_1(sb, 20u, 123u); - tint_atomicSub_1(sb, 20u, 123u); - tint_atomicMax_1(sb, 20u, 123u); - tint_atomicMin_1(sb, 20u, 123u); - tint_atomicAnd_1(sb, 20u, 123u); - tint_atomicOr_1(sb, 20u, 123u); - tint_atomicXor_1(sb, 20u, 123u); - tint_atomicExchange_1(sb, 20u, 123u); - tint_atomicCompareExchangeWeak_1(sb, 20u, 123u, 345u); + tint_atomicStore(&(sb), 16u, 123); + tint_atomicLoad(&(sb), 16u); + tint_atomicAdd(&(sb), 16u, 123); + tint_atomicSub(&(sb), 16u, 123); + tint_atomicMax(&(sb), 16u, 123); + tint_atomicMin(&(sb), 16u, 123); + tint_atomicAnd(&(sb), 16u, 123); + tint_atomicOr(&(sb), 16u, 123); + tint_atomicXor(&(sb), 16u, 123); + tint_atomicExchange(&(sb), 16u, 123); + tint_atomicCompareExchangeWeak(&(sb), 16u, 123, 345); + tint_atomicStore_1(&(sb), 20u, 123u); + tint_atomicLoad_1(&(sb), 20u); + tint_atomicAdd_1(&(sb), 20u, 123u); + tint_atomicSub_1(&(sb), 20u, 123u); + tint_atomicMax_1(&(sb), 20u, 123u); + tint_atomicMin_1(&(sb), 20u, 123u); + tint_atomicAnd_1(&(sb), 20u, 123u); + tint_atomicOr_1(&(sb), 20u, 123u); + tint_atomicXor_1(&(sb), 20u, 123u); + tint_atomicExchange_1(&(sb), 20u, 123u); + tint_atomicCompareExchangeWeak_1(&(sb), 20u, 123u, 345u); } )"; @@ -2614,34 +2614,34 @@ struct SB { auto* expect = R"( @internal(intrinsic_atomic_store_storage_i32) @internal(disable_validation__function_has_no_body) -fn tint_atomicStore(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) +fn tint_atomicStore(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, param_1 : i32) @internal(intrinsic_atomic_load_storage_i32) @internal(disable_validation__function_has_no_body) -fn tint_atomicLoad(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> i32 +fn tint_atomicLoad(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32) -> i32 @internal(intrinsic_atomic_add_storage_i32) @internal(disable_validation__function_has_no_body) -fn tint_atomicAdd(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32 +fn tint_atomicAdd(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, param_1 : i32) -> i32 @internal(intrinsic_atomic_sub_storage_i32) @internal(disable_validation__function_has_no_body) -fn tint_atomicSub(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32 +fn tint_atomicSub(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, param_1 : i32) -> i32 @internal(intrinsic_atomic_max_storage_i32) @internal(disable_validation__function_has_no_body) -fn tint_atomicMax(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32 +fn tint_atomicMax(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, param_1 : i32) -> i32 @internal(intrinsic_atomic_min_storage_i32) @internal(disable_validation__function_has_no_body) -fn tint_atomicMin(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32 +fn tint_atomicMin(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, param_1 : i32) -> i32 @internal(intrinsic_atomic_and_storage_i32) @internal(disable_validation__function_has_no_body) -fn tint_atomicAnd(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32 +fn tint_atomicAnd(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, param_1 : i32) -> i32 @internal(intrinsic_atomic_or_storage_i32) @internal(disable_validation__function_has_no_body) -fn tint_atomicOr(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32 +fn tint_atomicOr(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, param_1 : i32) -> i32 @internal(intrinsic_atomic_xor_storage_i32) @internal(disable_validation__function_has_no_body) -fn tint_atomicXor(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32 +fn tint_atomicXor(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, param_1 : i32) -> i32 @internal(intrinsic_atomic_exchange_storage_i32) @internal(disable_validation__function_has_no_body) -fn tint_atomicExchange(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32 +fn tint_atomicExchange(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, param_1 : i32) -> i32 struct atomic_compare_exchange_weak_ret_type { old_value : i32, @@ -2649,37 +2649,37 @@ struct atomic_compare_exchange_weak_ret_type { } @internal(intrinsic_atomic_compare_exchange_weak_storage_i32) @internal(disable_validation__function_has_no_body) -fn tint_atomicCompareExchangeWeak(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32, param_2 : i32) -> atomic_compare_exchange_weak_ret_type +fn tint_atomicCompareExchangeWeak(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, param_1 : i32, param_2 : i32) -> atomic_compare_exchange_weak_ret_type @internal(intrinsic_atomic_store_storage_u32) @internal(disable_validation__function_has_no_body) -fn tint_atomicStore_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) +fn tint_atomicStore_1(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, param_1 : u32) @internal(intrinsic_atomic_load_storage_u32) @internal(disable_validation__function_has_no_body) -fn tint_atomicLoad_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> u32 +fn tint_atomicLoad_1(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32) -> u32 @internal(intrinsic_atomic_add_storage_u32) @internal(disable_validation__function_has_no_body) -fn tint_atomicAdd_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32 +fn tint_atomicAdd_1(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, param_1 : u32) -> u32 @internal(intrinsic_atomic_sub_storage_u32) @internal(disable_validation__function_has_no_body) -fn tint_atomicSub_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32 +fn tint_atomicSub_1(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, param_1 : u32) -> u32 @internal(intrinsic_atomic_max_storage_u32) @internal(disable_validation__function_has_no_body) -fn tint_atomicMax_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32 +fn tint_atomicMax_1(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, param_1 : u32) -> u32 @internal(intrinsic_atomic_min_storage_u32) @internal(disable_validation__function_has_no_body) -fn tint_atomicMin_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32 +fn tint_atomicMin_1(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, param_1 : u32) -> u32 @internal(intrinsic_atomic_and_storage_u32) @internal(disable_validation__function_has_no_body) -fn tint_atomicAnd_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32 +fn tint_atomicAnd_1(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, param_1 : u32) -> u32 @internal(intrinsic_atomic_or_storage_u32) @internal(disable_validation__function_has_no_body) -fn tint_atomicOr_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32 +fn tint_atomicOr_1(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, param_1 : u32) -> u32 @internal(intrinsic_atomic_xor_storage_u32) @internal(disable_validation__function_has_no_body) -fn tint_atomicXor_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32 +fn tint_atomicXor_1(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, param_1 : u32) -> u32 @internal(intrinsic_atomic_exchange_storage_u32) @internal(disable_validation__function_has_no_body) -fn tint_atomicExchange_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32 +fn tint_atomicExchange_1(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, param_1 : u32) -> u32 struct atomic_compare_exchange_weak_ret_type_1 { old_value : u32, @@ -2687,32 +2687,32 @@ struct atomic_compare_exchange_weak_ret_type_1 { } @internal(intrinsic_atomic_compare_exchange_weak_storage_u32) @internal(disable_validation__function_has_no_body) -fn tint_atomicCompareExchangeWeak_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32, param_2 : u32) -> atomic_compare_exchange_weak_ret_type_1 +fn tint_atomicCompareExchangeWeak_1(@internal(disable_validation__function_parameter) buffer : ptr, offset : u32, param_1 : u32, param_2 : u32) -> atomic_compare_exchange_weak_ret_type_1 @compute @workgroup_size(1) fn main() { - tint_atomicStore(sb, 16u, 123); - tint_atomicLoad(sb, 16u); - tint_atomicAdd(sb, 16u, 123); - tint_atomicSub(sb, 16u, 123); - tint_atomicMax(sb, 16u, 123); - tint_atomicMin(sb, 16u, 123); - tint_atomicAnd(sb, 16u, 123); - tint_atomicOr(sb, 16u, 123); - tint_atomicXor(sb, 16u, 123); - tint_atomicExchange(sb, 16u, 123); - tint_atomicCompareExchangeWeak(sb, 16u, 123, 345); - tint_atomicStore_1(sb, 20u, 123u); - tint_atomicLoad_1(sb, 20u); - tint_atomicAdd_1(sb, 20u, 123u); - tint_atomicSub_1(sb, 20u, 123u); - tint_atomicMax_1(sb, 20u, 123u); - tint_atomicMin_1(sb, 20u, 123u); - tint_atomicAnd_1(sb, 20u, 123u); - tint_atomicOr_1(sb, 20u, 123u); - tint_atomicXor_1(sb, 20u, 123u); - tint_atomicExchange_1(sb, 20u, 123u); - tint_atomicCompareExchangeWeak_1(sb, 20u, 123u, 345u); + tint_atomicStore(&(sb), 16u, 123); + tint_atomicLoad(&(sb), 16u); + tint_atomicAdd(&(sb), 16u, 123); + tint_atomicSub(&(sb), 16u, 123); + tint_atomicMax(&(sb), 16u, 123); + tint_atomicMin(&(sb), 16u, 123); + tint_atomicAnd(&(sb), 16u, 123); + tint_atomicOr(&(sb), 16u, 123); + tint_atomicXor(&(sb), 16u, 123); + tint_atomicExchange(&(sb), 16u, 123); + tint_atomicCompareExchangeWeak(&(sb), 16u, 123, 345); + tint_atomicStore_1(&(sb), 20u, 123u); + tint_atomicLoad_1(&(sb), 20u); + tint_atomicAdd_1(&(sb), 20u, 123u); + tint_atomicSub_1(&(sb), 20u, 123u); + tint_atomicMax_1(&(sb), 20u, 123u); + tint_atomicMin_1(&(sb), 20u, 123u); + tint_atomicAnd_1(&(sb), 20u, 123u); + tint_atomicOr_1(&(sb), 20u, 123u); + tint_atomicXor_1(&(sb), 20u, 123u); + tint_atomicExchange_1(&(sb), 20u, 123u); + tint_atomicCompareExchangeWeak_1(&(sb), 20u, 123u, 345u); } @group(0) @binding(0) var sb : SB; diff --git a/src/tint/transform/manager.cc b/src/tint/transform/manager.cc index e5f7682ead..4e83320b41 100644 --- a/src/tint/transform/manager.cc +++ b/src/tint/transform/manager.cc @@ -49,7 +49,8 @@ Output Manager::Run(const Program* program, const DataMap& data) const { Output out; for (const auto& transform : transforms_) { if (!transform->ShouldRun(in, data)) { - TINT_IF_PRINT_PROGRAM(std::cout << "Skipping " << transform->TypeInfo().name << std::endl); + TINT_IF_PRINT_PROGRAM(std::cout << "Skipping " << transform->TypeInfo().name + << std::endl); continue; } TINT_IF_PRINT_PROGRAM(print_program("Input to", transform.get())); diff --git a/src/tint/writer/hlsl/generator_impl.cc b/src/tint/writer/hlsl/generator_impl.cc index 19af4fad52..90b50ac529 100644 --- a/src/tint/writer/hlsl/generator_impl.cc +++ b/src/tint/writer/hlsl/generator_impl.cc @@ -2776,14 +2776,25 @@ bool GeneratorImpl::EmitFunction(const ast::Function* func) { first = false; auto const* type = v->Type(); + auto storage_class = ast::StorageClass::kNone; + auto access = ast::Access::kUndefined; if (auto* ptr = type->As()) { - // Transform pointer parameters in to `inout` parameters. - // The WGSL spec is highly restrictive in what can be passed in pointer - // parameters, which allows for this transformation. See: - // https://gpuweb.github.io/gpuweb/wgsl/#function-restriction - out << "inout "; type = ptr->StoreType(); + switch (ptr->StorageClass()) { + case ast::StorageClass::kStorage: + case ast::StorageClass::kUniform: + // Not allowed by WGSL, but is used by certain transforms (e.g. DMA) to pass + // storage buffers and uniform buffers down into transform-generated + // functions. In this situation we want to generate the parameter without an + // 'inout', using the storage class and access from the pointer. + storage_class = ptr->StorageClass(); + access = ptr->Access(); + break; + default: + // Transform regular WGSL pointer parameters in to `inout` parameters. + out << "inout "; + } } // Note: WGSL only allows for StorageClass::kNone on parameters, however @@ -2792,7 +2803,7 @@ bool GeneratorImpl::EmitFunction(const ast::Function* func) { // StorageClass::kStorage or StorageClass::kUniform. This is required to // correctly translate the parameter to a [RW]ByteAddressBuffer for // storage buffers and a uint4[N] for uniform buffers. - if (!EmitTypeAndName(out, type, v->StorageClass(), v->Access(), + if (!EmitTypeAndName(out, type, storage_class, access, builder_.Symbols().NameFor(v->Declaration()->symbol))) { return false; }