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 <bclayton@google.com>
Reviewed-by: Dan Sinclair <dsinclair@chromium.org>
diff --git a/src/tint/ast/disable_validation_attribute.cc b/src/tint/ast/disable_validation_attribute.cc
index 4bc9f74..c5a6545 100644
--- a/src/tint/ast/disable_validation_attribute.cc
+++ b/src/tint/ast/disable_validation_attribute.cc
@@ -35,8 +35,8 @@
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 db70ad4..e44f7b8 100644
--- a/src/tint/ast/disable_validation_attribute.h
+++ b/src/tint/ast/disable_validation_attribute.h
@@ -24,27 +24,23 @@
/// 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 2b4d8a9..4d1ee07 100644
--- a/src/tint/resolver/validator.cc
+++ b/src/tint/resolver/validator.cc
@@ -722,19 +722,20 @@
auto* decl = var->Declaration();
+ if (IsValidationDisabled(decl->attributes, ast::DisabledValidation::kFunctionParameter)) {
+ return true;
+ }
+
for (auto* attr : decl->attributes) {
if (!func->IsEntryPoint() && !attr->Is<ast::InternalAttribute>()) {
AddError("attribute is not valid for non-entry point function parameters",
attr->source);
return false;
- } else if (!attr->IsAnyOf<ast::BuiltinAttribute, ast::InvariantAttribute,
- ast::LocationAttribute, ast::InterpolateAttribute,
- ast::InternalAttribute>() &&
- (IsValidationEnabled(decl->attributes,
- ast::DisabledValidation::kEntryPointParameter) &&
- IsValidationEnabled(
- decl->attributes,
- ast::DisabledValidation::kIgnoreConstructibleFunctionParameter))) {
+ }
+ if (!attr->IsAnyOf<ast::BuiltinAttribute, ast::InvariantAttribute, ast::LocationAttribute,
+ ast::InterpolateAttribute, ast::InternalAttribute>() &&
+ (IsValidationEnabled(decl->attributes,
+ ast::DisabledValidation::kEntryPointParameter))) {
AddError("attribute is not valid for function parameters", attr->source);
return false;
}
@@ -753,9 +754,7 @@
}
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 @@
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 9cf1175..acf55a6 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 @@
// 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<const sem::Type*, Symbol> buffer_size_intrinsics;
- auto get_buffer_size_intrinsic = [&](const sem::Type* buffer_type) {
+ std::unordered_map<const sem::Reference*, Symbol> 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<ast::Function>(
name,
ast::ParameterList{
- // Note: The buffer parameter requires the kStorage StorageClass
- // in order for HLSL to emit this as a ByteAddressBuffer.
- ctx.dst->create<ast::Variable>(ctx.dst->Sym("buffer"),
- ast::StorageClass::kStorage,
- 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 @@
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 @@
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<sem::Reference>();
- // 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 @@
// 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 @@
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<sem::Struct>()) {
- // 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<sem::Array>();
- total_size =
- ctx.dst->Sub(total_size, u32(array_member_sem->Offset()));
- } else if (auto* arr = storage_buffer_type->As<sem::Array>()) {
- 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<sem::Array>();
+ },
+ [&](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 e2674b0..98f2ed7 100644
--- a/src/tint/transform/calculate_array_length_test.cc
+++ b/src/tint/transform/calculate_array_length_test.cc
@@ -76,14 +76,14 @@
auto* expect = R"(
@internal(intrinsic_buffer_size)
-fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : array<i32>, result : ptr<function, u32>)
+fn tint_symbol(@internal(disable_validation__function_parameter) buffer : ptr<storage, array<i32>, read>, result : ptr<function, u32>)
@group(0) @binding(0) var<storage, read> sb : array<i32>;
@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 @@
auto* expect = R"(
@internal(intrinsic_buffer_size)
-fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, result : ptr<function, u32>)
+fn tint_symbol(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read>, result : ptr<function, u32>)
struct SB {
x : i32,
@@ -123,7 +123,7 @@
@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 @@
)";
auto* expect = R"(
@internal(intrinsic_buffer_size)
-fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : array<S>, result : ptr<function, u32>)
+fn tint_symbol(@internal(disable_validation__function_parameter) buffer : ptr<storage, array<S>, read>, result : ptr<function, u32>)
struct S {
f : f32,
@@ -160,7 +160,7 @@
@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 @@
)";
auto* expect = R"(
@internal(intrinsic_buffer_size)
-fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : array<array<S, 4u>>, result : ptr<function, u32>)
+fn tint_symbol(@internal(disable_validation__function_parameter) buffer : ptr<storage, array<array<S, 4u>>, read>, result : ptr<function, u32>)
struct S {
f : f32,
@@ -197,7 +197,7 @@
@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 @@
auto* expect = R"(
@internal(intrinsic_buffer_size)
-fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : array<i32>, result : ptr<function, u32>)
+fn tint_symbol(@internal(disable_validation__function_parameter) buffer : ptr<storage, array<i32>, read>, result : ptr<function, u32>)
@group(0) @binding(0) var<storage, read> sb : array<i32>;
@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 @@
auto* expect = R"(
@internal(intrinsic_buffer_size)
-fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, result : ptr<function, u32>)
+fn tint_symbol(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read>, result : ptr<function, u32>)
struct SB {
x : i32,
@@ -273,7 +273,7 @@
@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 @@
auto* expect = R"(
@internal(intrinsic_buffer_size)
-fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, result : ptr<function, u32>)
+fn tint_symbol(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read>, result : ptr<function, u32>)
struct SB {
x : i32,
@@ -322,13 +322,13 @@
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 @@
auto* expect = R"(
@internal(intrinsic_buffer_size)
-fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB1, result : ptr<function, u32>)
+fn tint_symbol(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB1, read>, result : ptr<function, u32>)
@internal(intrinsic_buffer_size)
-fn tint_symbol_3(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB2, result : ptr<function, u32>)
+fn tint_symbol_3(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB2, read>, result : ptr<function, u32>)
@internal(intrinsic_buffer_size)
-fn tint_symbol_6(@internal(disable_validation__ignore_constructible_function_parameter) buffer : array<i32>, result : ptr<function, u32>)
+fn tint_symbol_6(@internal(disable_validation__function_parameter) buffer : ptr<storage, array<i32>, read>, result : ptr<function, u32>)
struct SB1 {
x : i32,
@@ -397,13 +397,13 @@
@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 @@
auto* expect =
R"(
@internal(intrinsic_buffer_size)
-fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, result : ptr<function, u32>)
+fn tint_symbol(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read>, result : ptr<function, u32>)
struct SB {
x : i32,
@@ -454,12 +454,12 @@
@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 @@
auto* expect = R"(
@internal(intrinsic_buffer_size)
-fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB1, result : ptr<function, u32>)
+fn tint_symbol(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB1, read>, result : ptr<function, u32>)
@internal(intrinsic_buffer_size)
-fn tint_symbol_3(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB2, result : ptr<function, u32>)
+fn tint_symbol_3(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB2, read>, result : ptr<function, u32>)
@internal(intrinsic_buffer_size)
-fn tint_symbol_6(@internal(disable_validation__ignore_constructible_function_parameter) buffer : array<i32>, result : ptr<function, u32>)
+fn tint_symbol_6(@internal(disable_validation__function_parameter) buffer : ptr<storage, array<i32>, read>, result : ptr<function, u32>)
@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 f89d9ab..32ad153 100644
--- a/src/tint/transform/decompose_memory_access.cc
+++ b/src/tint/transform/decompose_memory_access.cc
@@ -98,29 +98,32 @@
/// 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 @@
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<SC, buf_ty, A>, 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 @@
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<ast::Variable>(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<ast::Function>(
+ 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<sem::Array>()) {
+ // fn load_func(buffer : buf_ty, offset : u32) -> array<T, N> {
+ // var arr : array<T, N>;
+ // for (var i = 0u; i < array_count; i = i + 1) {
+ // arr[i] = el_load_func(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::BinaryExpression>(
+ 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<ast::Function>(
- 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<sem::Array>()) {
- // fn load_func(buf : buf_ty, offset : u32) -> array<T, N> {
- // var arr : array<T, N>;
- // for (var i = 0u; i < array_count; i = i + 1) {
- // arr[i] = el_load_func(buf, offset + i * array_stride)
- // }
- // return arr;
- // }
- auto load = LoadFunc(buf_ty, arr_ty->ElemType()->UnwrapRef(), var_user);
- auto* arr = b.Var(b.Symbols().New("arr"), CreateASTTypeFor(ctx, arr_ty));
- auto* i = b.Var(b.Symbols().New("i"), nullptr, b.Expr(0_u));
- auto* for_init = b.Decl(i);
- auto* for_cond = b.create<ast::BinaryExpression>(
- 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<sem::Matrix>()) {
- 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));
+ b.Func(name, params, CreateASTTypeFor(ctx, arr_ty),
+ {
+ b.Decl(arr),
+ for_loop,
+ b.Return(arr),
+ });
+ } else {
+ ast::ExpressionList values;
+ if (auto* mat_ty = el_ty->As<sem::Matrix>()) {
+ auto* vec_ty = mat_ty->ColumnType();
+ Symbol load = LoadFunc(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<sem::Struct>()) {
+ 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));
+ }
}
- } else if (auto* str = el_ty->As<sem::Struct>()) {
- 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<SC, buf_ty, A>, 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 @@
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<ast::Variable>(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<ast::Function>(
+ 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::BinaryExpression>(
+ 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<ast::Function>(
- 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<sem::Array>()) {
- // fn store_func(buf : buf_ty, offset : u32, value : el_ty) {
- // var array = value; // No dynamic indexing on constant arrays
- // for (var i = 0u; i < array_count; i = i + 1) {
- // arr[i] = el_store_func(buf, offset + i * array_stride,
- // value[i])
- // }
- // return arr;
- // }
- auto* array = b.Var(b.Symbols().New("array"), nullptr, b.Expr("value"));
- auto store = StoreFunc(buf_ty, arr_ty->ElemType()->UnwrapRef(), var_user);
- auto* i = b.Var(b.Symbols().New("i"), nullptr, b.Expr(0_u));
- auto* for_init = b.Decl(i);
- auto* for_cond = b.create<ast::BinaryExpression>(
- 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<sem::Matrix>()) {
- 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<sem::Struct>()) {
- 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<storage, buf_ty, A>, 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 @@
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<ast::Variable>(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 @@
if (auto* builtin = call->Target()->As<sem::Builtin>()) {
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 @@
Symbol func = state.AtomicFunc(buf_ty, el_ty, builtin,
access.var->As<sem::VariableUser>());
- 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 @@
}
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<sem::VariableUser>());
- 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<sem::VariableUser>());
- 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 4b96bcb..581731e 100644
--- a/src/tint/transform/decompose_memory_access_test.cc
+++ b/src/tint/transform/decompose_memory_access_test.cc
@@ -134,78 +134,78 @@
@group(0) @binding(0) var<storage, read_write> sb : SB;
@internal(intrinsic_load_storage_i32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> i32
+fn tint_symbol(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, 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<storage, SB, read_write>, 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<storage, SB, read_write>, offset : u32) -> f32
@internal(intrinsic_load_storage_vec2_i32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_3(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> vec2<i32>
+fn tint_symbol_3(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> vec2<i32>
@internal(intrinsic_load_storage_vec2_u32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_4(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> vec2<u32>
+fn tint_symbol_4(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> vec2<u32>
@internal(intrinsic_load_storage_vec2_f32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_5(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> vec2<f32>
+fn tint_symbol_5(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> vec2<f32>
@internal(intrinsic_load_storage_vec3_i32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_6(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> vec3<i32>
+fn tint_symbol_6(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> vec3<i32>
@internal(intrinsic_load_storage_vec3_u32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_7(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> vec3<u32>
+fn tint_symbol_7(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> vec3<u32>
@internal(intrinsic_load_storage_vec3_f32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_8(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> vec3<f32>
+fn tint_symbol_8(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> vec3<f32>
@internal(intrinsic_load_storage_vec4_i32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_9(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> vec4<i32>
+fn tint_symbol_9(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> vec4<i32>
@internal(intrinsic_load_storage_vec4_u32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_10(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> vec4<u32>
+fn tint_symbol_10(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> vec4<u32>
@internal(intrinsic_load_storage_vec4_f32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_11(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> vec4<f32>
+fn tint_symbol_11(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> vec4<f32>
-fn tint_symbol_12(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> mat2x2<f32> {
+fn tint_symbol_12(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> mat2x2<f32> {
return mat2x2<f32>(tint_symbol_5(buffer, (offset + 0u)), tint_symbol_5(buffer, (offset + 8u)));
}
-fn tint_symbol_13(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> mat2x3<f32> {
+fn tint_symbol_13(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> mat2x3<f32> {
return mat2x3<f32>(tint_symbol_8(buffer, (offset + 0u)), tint_symbol_8(buffer, (offset + 16u)));
}
-fn tint_symbol_14(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> mat2x4<f32> {
+fn tint_symbol_14(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> mat2x4<f32> {
return mat2x4<f32>(tint_symbol_11(buffer, (offset + 0u)), tint_symbol_11(buffer, (offset + 16u)));
}
-fn tint_symbol_15(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> mat3x2<f32> {
+fn tint_symbol_15(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> mat3x2<f32> {
return mat3x2<f32>(tint_symbol_5(buffer, (offset + 0u)), tint_symbol_5(buffer, (offset + 8u)), tint_symbol_5(buffer, (offset + 16u)));
}
-fn tint_symbol_16(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> mat3x3<f32> {
+fn tint_symbol_16(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> mat3x3<f32> {
return mat3x3<f32>(tint_symbol_8(buffer, (offset + 0u)), tint_symbol_8(buffer, (offset + 16u)), tint_symbol_8(buffer, (offset + 32u)));
}
-fn tint_symbol_17(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> mat3x4<f32> {
+fn tint_symbol_17(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> mat3x4<f32> {
return mat3x4<f32>(tint_symbol_11(buffer, (offset + 0u)), tint_symbol_11(buffer, (offset + 16u)), tint_symbol_11(buffer, (offset + 32u)));
}
-fn tint_symbol_18(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> mat4x2<f32> {
+fn tint_symbol_18(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> mat4x2<f32> {
return mat4x2<f32>(tint_symbol_5(buffer, (offset + 0u)), tint_symbol_5(buffer, (offset + 8u)), tint_symbol_5(buffer, (offset + 16u)), tint_symbol_5(buffer, (offset + 24u)));
}
-fn tint_symbol_19(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> mat4x3<f32> {
+fn tint_symbol_19(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> mat4x3<f32> {
return mat4x3<f32>(tint_symbol_8(buffer, (offset + 0u)), tint_symbol_8(buffer, (offset + 16u)), tint_symbol_8(buffer, (offset + 32u)), tint_symbol_8(buffer, (offset + 48u)));
}
-fn tint_symbol_20(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> mat4x4<f32> {
+fn tint_symbol_20(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> mat4x4<f32> {
return mat4x4<f32>(tint_symbol_11(buffer, (offset + 0u)), tint_symbol_11(buffer, (offset + 16u)), tint_symbol_11(buffer, (offset + 32u)), tint_symbol_11(buffer, (offset + 48u)));
}
-fn tint_symbol_21(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> array<vec3<f32>, 2u> {
+fn tint_symbol_21(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> array<vec3<f32>, 2u> {
var arr : array<vec3<f32>, 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 @@
@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<i32> = tint_symbol_3(sb, 16u);
- var e : vec2<u32> = tint_symbol_4(sb, 24u);
- var f : vec2<f32> = tint_symbol_5(sb, 32u);
- var g : vec3<i32> = tint_symbol_6(sb, 48u);
- var h : vec3<u32> = tint_symbol_7(sb, 64u);
- var i : vec3<f32> = tint_symbol_8(sb, 80u);
- var j : vec4<i32> = tint_symbol_9(sb, 96u);
- var k : vec4<u32> = tint_symbol_10(sb, 112u);
- var l : vec4<f32> = tint_symbol_11(sb, 128u);
- var m : mat2x2<f32> = tint_symbol_12(sb, 144u);
- var n : mat2x3<f32> = tint_symbol_13(sb, 160u);
- var o : mat2x4<f32> = tint_symbol_14(sb, 192u);
- var p : mat3x2<f32> = tint_symbol_15(sb, 224u);
- var q : mat3x3<f32> = tint_symbol_16(sb, 256u);
- var r : mat3x4<f32> = tint_symbol_17(sb, 304u);
- var s : mat4x2<f32> = tint_symbol_18(sb, 352u);
- var t : mat4x3<f32> = tint_symbol_19(sb, 384u);
- var u : mat4x4<f32> = tint_symbol_20(sb, 448u);
- var v : array<vec3<f32>, 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<i32> = tint_symbol_3(&(sb), 16u);
+ var e : vec2<u32> = tint_symbol_4(&(sb), 24u);
+ var f : vec2<f32> = tint_symbol_5(&(sb), 32u);
+ var g : vec3<i32> = tint_symbol_6(&(sb), 48u);
+ var h : vec3<u32> = tint_symbol_7(&(sb), 64u);
+ var i : vec3<f32> = tint_symbol_8(&(sb), 80u);
+ var j : vec4<i32> = tint_symbol_9(&(sb), 96u);
+ var k : vec4<u32> = tint_symbol_10(&(sb), 112u);
+ var l : vec4<f32> = tint_symbol_11(&(sb), 128u);
+ var m : mat2x2<f32> = tint_symbol_12(&(sb), 144u);
+ var n : mat2x3<f32> = tint_symbol_13(&(sb), 160u);
+ var o : mat2x4<f32> = tint_symbol_14(&(sb), 192u);
+ var p : mat3x2<f32> = tint_symbol_15(&(sb), 224u);
+ var q : mat3x3<f32> = tint_symbol_16(&(sb), 256u);
+ var r : mat3x4<f32> = tint_symbol_17(&(sb), 304u);
+ var s : mat4x2<f32> = tint_symbol_18(&(sb), 352u);
+ var t : mat4x3<f32> = tint_symbol_19(&(sb), 384u);
+ var u : mat4x4<f32> = tint_symbol_20(&(sb), 448u);
+ var v : array<vec3<f32>, 2> = tint_symbol_21(&(sb), 512u);
}
)";
@@ -303,78 +303,78 @@
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<storage, SB, read_write>, 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<storage, SB, read_write>, 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<storage, SB, read_write>, offset : u32) -> f32
@internal(intrinsic_load_storage_vec2_i32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_3(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> vec2<i32>
+fn tint_symbol_3(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> vec2<i32>
@internal(intrinsic_load_storage_vec2_u32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_4(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> vec2<u32>
+fn tint_symbol_4(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> vec2<u32>
@internal(intrinsic_load_storage_vec2_f32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_5(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> vec2<f32>
+fn tint_symbol_5(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> vec2<f32>
@internal(intrinsic_load_storage_vec3_i32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_6(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> vec3<i32>
+fn tint_symbol_6(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> vec3<i32>
@internal(intrinsic_load_storage_vec3_u32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_7(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> vec3<u32>
+fn tint_symbol_7(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> vec3<u32>
@internal(intrinsic_load_storage_vec3_f32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_8(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> vec3<f32>
+fn tint_symbol_8(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> vec3<f32>
@internal(intrinsic_load_storage_vec4_i32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_9(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> vec4<i32>
+fn tint_symbol_9(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> vec4<i32>
@internal(intrinsic_load_storage_vec4_u32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_10(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> vec4<u32>
+fn tint_symbol_10(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> vec4<u32>
@internal(intrinsic_load_storage_vec4_f32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_11(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> vec4<f32>
+fn tint_symbol_11(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> vec4<f32>
-fn tint_symbol_12(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> mat2x2<f32> {
+fn tint_symbol_12(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> mat2x2<f32> {
return mat2x2<f32>(tint_symbol_5(buffer, (offset + 0u)), tint_symbol_5(buffer, (offset + 8u)));
}
-fn tint_symbol_13(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> mat2x3<f32> {
+fn tint_symbol_13(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> mat2x3<f32> {
return mat2x3<f32>(tint_symbol_8(buffer, (offset + 0u)), tint_symbol_8(buffer, (offset + 16u)));
}
-fn tint_symbol_14(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> mat2x4<f32> {
+fn tint_symbol_14(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> mat2x4<f32> {
return mat2x4<f32>(tint_symbol_11(buffer, (offset + 0u)), tint_symbol_11(buffer, (offset + 16u)));
}
-fn tint_symbol_15(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> mat3x2<f32> {
+fn tint_symbol_15(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> mat3x2<f32> {
return mat3x2<f32>(tint_symbol_5(buffer, (offset + 0u)), tint_symbol_5(buffer, (offset + 8u)), tint_symbol_5(buffer, (offset + 16u)));
}
-fn tint_symbol_16(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> mat3x3<f32> {
+fn tint_symbol_16(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> mat3x3<f32> {
return mat3x3<f32>(tint_symbol_8(buffer, (offset + 0u)), tint_symbol_8(buffer, (offset + 16u)), tint_symbol_8(buffer, (offset + 32u)));
}
-fn tint_symbol_17(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> mat3x4<f32> {
+fn tint_symbol_17(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> mat3x4<f32> {
return mat3x4<f32>(tint_symbol_11(buffer, (offset + 0u)), tint_symbol_11(buffer, (offset + 16u)), tint_symbol_11(buffer, (offset + 32u)));
}
-fn tint_symbol_18(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> mat4x2<f32> {
+fn tint_symbol_18(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> mat4x2<f32> {
return mat4x2<f32>(tint_symbol_5(buffer, (offset + 0u)), tint_symbol_5(buffer, (offset + 8u)), tint_symbol_5(buffer, (offset + 16u)), tint_symbol_5(buffer, (offset + 24u)));
}
-fn tint_symbol_19(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> mat4x3<f32> {
+fn tint_symbol_19(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> mat4x3<f32> {
return mat4x3<f32>(tint_symbol_8(buffer, (offset + 0u)), tint_symbol_8(buffer, (offset + 16u)), tint_symbol_8(buffer, (offset + 32u)), tint_symbol_8(buffer, (offset + 48u)));
}
-fn tint_symbol_20(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> mat4x4<f32> {
+fn tint_symbol_20(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> mat4x4<f32> {
return mat4x4<f32>(tint_symbol_11(buffer, (offset + 0u)), tint_symbol_11(buffer, (offset + 16u)), tint_symbol_11(buffer, (offset + 32u)), tint_symbol_11(buffer, (offset + 48u)));
}
-fn tint_symbol_21(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> array<vec3<f32>, 2u> {
+fn tint_symbol_21(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> array<vec3<f32>, 2u> {
var arr : array<vec3<f32>, 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 @@
@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<i32> = tint_symbol_3(sb, 16u);
- var e : vec2<u32> = tint_symbol_4(sb, 24u);
- var f : vec2<f32> = tint_symbol_5(sb, 32u);
- var g : vec3<i32> = tint_symbol_6(sb, 48u);
- var h : vec3<u32> = tint_symbol_7(sb, 64u);
- var i : vec3<f32> = tint_symbol_8(sb, 80u);
- var j : vec4<i32> = tint_symbol_9(sb, 96u);
- var k : vec4<u32> = tint_symbol_10(sb, 112u);
- var l : vec4<f32> = tint_symbol_11(sb, 128u);
- var m : mat2x2<f32> = tint_symbol_12(sb, 144u);
- var n : mat2x3<f32> = tint_symbol_13(sb, 160u);
- var o : mat2x4<f32> = tint_symbol_14(sb, 192u);
- var p : mat3x2<f32> = tint_symbol_15(sb, 224u);
- var q : mat3x3<f32> = tint_symbol_16(sb, 256u);
- var r : mat3x4<f32> = tint_symbol_17(sb, 304u);
- var s : mat4x2<f32> = tint_symbol_18(sb, 352u);
- var t : mat4x3<f32> = tint_symbol_19(sb, 384u);
- var u : mat4x4<f32> = tint_symbol_20(sb, 448u);
- var v : array<vec3<f32>, 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<i32> = tint_symbol_3(&(sb), 16u);
+ var e : vec2<u32> = tint_symbol_4(&(sb), 24u);
+ var f : vec2<f32> = tint_symbol_5(&(sb), 32u);
+ var g : vec3<i32> = tint_symbol_6(&(sb), 48u);
+ var h : vec3<u32> = tint_symbol_7(&(sb), 64u);
+ var i : vec3<f32> = tint_symbol_8(&(sb), 80u);
+ var j : vec4<i32> = tint_symbol_9(&(sb), 96u);
+ var k : vec4<u32> = tint_symbol_10(&(sb), 112u);
+ var l : vec4<f32> = tint_symbol_11(&(sb), 128u);
+ var m : mat2x2<f32> = tint_symbol_12(&(sb), 144u);
+ var n : mat2x3<f32> = tint_symbol_13(&(sb), 160u);
+ var o : mat2x4<f32> = tint_symbol_14(&(sb), 192u);
+ var p : mat3x2<f32> = tint_symbol_15(&(sb), 224u);
+ var q : mat3x3<f32> = tint_symbol_16(&(sb), 256u);
+ var r : mat3x4<f32> = tint_symbol_17(&(sb), 304u);
+ var s : mat4x2<f32> = tint_symbol_18(&(sb), 352u);
+ var t : mat4x3<f32> = tint_symbol_19(&(sb), 384u);
+ var u : mat4x4<f32> = tint_symbol_20(&(sb), 448u);
+ var v : array<vec3<f32>, 2> = tint_symbol_21(&(sb), 512u);
}
@group(0) @binding(0) var<storage, read_write> sb : SB;
@@ -526,78 +526,78 @@
@group(0) @binding(0) var<uniform> 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<uniform, UB, read>, 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<uniform, UB, read>, 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<uniform, UB, read>, 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<i32>
+fn tint_symbol_3(@internal(disable_validation__function_parameter) buffer : ptr<uniform, UB, read>, offset : u32) -> vec2<i32>
@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<u32>
+fn tint_symbol_4(@internal(disable_validation__function_parameter) buffer : ptr<uniform, UB, read>, offset : u32) -> vec2<u32>
@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<f32>
+fn tint_symbol_5(@internal(disable_validation__function_parameter) buffer : ptr<uniform, UB, read>, offset : u32) -> vec2<f32>
@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<i32>
+fn tint_symbol_6(@internal(disable_validation__function_parameter) buffer : ptr<uniform, UB, read>, offset : u32) -> vec3<i32>
@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<u32>
+fn tint_symbol_7(@internal(disable_validation__function_parameter) buffer : ptr<uniform, UB, read>, offset : u32) -> vec3<u32>
@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<f32>
+fn tint_symbol_8(@internal(disable_validation__function_parameter) buffer : ptr<uniform, UB, read>, offset : u32) -> vec3<f32>
@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<i32>
+fn tint_symbol_9(@internal(disable_validation__function_parameter) buffer : ptr<uniform, UB, read>, offset : u32) -> vec4<i32>
@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<u32>
+fn tint_symbol_10(@internal(disable_validation__function_parameter) buffer : ptr<uniform, UB, read>, offset : u32) -> vec4<u32>
@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<f32>
+fn tint_symbol_11(@internal(disable_validation__function_parameter) buffer : ptr<uniform, UB, read>, offset : u32) -> vec4<f32>
-fn tint_symbol_12(@internal(disable_validation__ignore_constructible_function_parameter) buffer : UB, offset : u32) -> mat2x2<f32> {
+fn tint_symbol_12(@internal(disable_validation__function_parameter) buffer : ptr<uniform, UB, read>, offset : u32) -> mat2x2<f32> {
return mat2x2<f32>(tint_symbol_5(buffer, (offset + 0u)), tint_symbol_5(buffer, (offset + 8u)));
}
-fn tint_symbol_13(@internal(disable_validation__ignore_constructible_function_parameter) buffer : UB, offset : u32) -> mat2x3<f32> {
+fn tint_symbol_13(@internal(disable_validation__function_parameter) buffer : ptr<uniform, UB, read>, offset : u32) -> mat2x3<f32> {
return mat2x3<f32>(tint_symbol_8(buffer, (offset + 0u)), tint_symbol_8(buffer, (offset + 16u)));
}
-fn tint_symbol_14(@internal(disable_validation__ignore_constructible_function_parameter) buffer : UB, offset : u32) -> mat2x4<f32> {
+fn tint_symbol_14(@internal(disable_validation__function_parameter) buffer : ptr<uniform, UB, read>, offset : u32) -> mat2x4<f32> {
return mat2x4<f32>(tint_symbol_11(buffer, (offset + 0u)), tint_symbol_11(buffer, (offset + 16u)));
}
-fn tint_symbol_15(@internal(disable_validation__ignore_constructible_function_parameter) buffer : UB, offset : u32) -> mat3x2<f32> {
+fn tint_symbol_15(@internal(disable_validation__function_parameter) buffer : ptr<uniform, UB, read>, offset : u32) -> mat3x2<f32> {
return mat3x2<f32>(tint_symbol_5(buffer, (offset + 0u)), tint_symbol_5(buffer, (offset + 8u)), tint_symbol_5(buffer, (offset + 16u)));
}
-fn tint_symbol_16(@internal(disable_validation__ignore_constructible_function_parameter) buffer : UB, offset : u32) -> mat3x3<f32> {
+fn tint_symbol_16(@internal(disable_validation__function_parameter) buffer : ptr<uniform, UB, read>, offset : u32) -> mat3x3<f32> {
return mat3x3<f32>(tint_symbol_8(buffer, (offset + 0u)), tint_symbol_8(buffer, (offset + 16u)), tint_symbol_8(buffer, (offset + 32u)));
}
-fn tint_symbol_17(@internal(disable_validation__ignore_constructible_function_parameter) buffer : UB, offset : u32) -> mat3x4<f32> {
+fn tint_symbol_17(@internal(disable_validation__function_parameter) buffer : ptr<uniform, UB, read>, offset : u32) -> mat3x4<f32> {
return mat3x4<f32>(tint_symbol_11(buffer, (offset + 0u)), tint_symbol_11(buffer, (offset + 16u)), tint_symbol_11(buffer, (offset + 32u)));
}
-fn tint_symbol_18(@internal(disable_validation__ignore_constructible_function_parameter) buffer : UB, offset : u32) -> mat4x2<f32> {
+fn tint_symbol_18(@internal(disable_validation__function_parameter) buffer : ptr<uniform, UB, read>, offset : u32) -> mat4x2<f32> {
return mat4x2<f32>(tint_symbol_5(buffer, (offset + 0u)), tint_symbol_5(buffer, (offset + 8u)), tint_symbol_5(buffer, (offset + 16u)), tint_symbol_5(buffer, (offset + 24u)));
}
-fn tint_symbol_19(@internal(disable_validation__ignore_constructible_function_parameter) buffer : UB, offset : u32) -> mat4x3<f32> {
+fn tint_symbol_19(@internal(disable_validation__function_parameter) buffer : ptr<uniform, UB, read>, offset : u32) -> mat4x3<f32> {
return mat4x3<f32>(tint_symbol_8(buffer, (offset + 0u)), tint_symbol_8(buffer, (offset + 16u)), tint_symbol_8(buffer, (offset + 32u)), tint_symbol_8(buffer, (offset + 48u)));
}
-fn tint_symbol_20(@internal(disable_validation__ignore_constructible_function_parameter) buffer : UB, offset : u32) -> mat4x4<f32> {
+fn tint_symbol_20(@internal(disable_validation__function_parameter) buffer : ptr<uniform, UB, read>, offset : u32) -> mat4x4<f32> {
return mat4x4<f32>(tint_symbol_11(buffer, (offset + 0u)), tint_symbol_11(buffer, (offset + 16u)), tint_symbol_11(buffer, (offset + 32u)), tint_symbol_11(buffer, (offset + 48u)));
}
-fn tint_symbol_21(@internal(disable_validation__ignore_constructible_function_parameter) buffer : UB, offset : u32) -> array<vec3<f32>, 2u> {
+fn tint_symbol_21(@internal(disable_validation__function_parameter) buffer : ptr<uniform, UB, read>, offset : u32) -> array<vec3<f32>, 2u> {
var arr : array<vec3<f32>, 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 @@
@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<i32> = tint_symbol_3(ub, 16u);
- var e : vec2<u32> = tint_symbol_4(ub, 24u);
- var f : vec2<f32> = tint_symbol_5(ub, 32u);
- var g : vec3<i32> = tint_symbol_6(ub, 48u);
- var h : vec3<u32> = tint_symbol_7(ub, 64u);
- var i : vec3<f32> = tint_symbol_8(ub, 80u);
- var j : vec4<i32> = tint_symbol_9(ub, 96u);
- var k : vec4<u32> = tint_symbol_10(ub, 112u);
- var l : vec4<f32> = tint_symbol_11(ub, 128u);
- var m : mat2x2<f32> = tint_symbol_12(ub, 144u);
- var n : mat2x3<f32> = tint_symbol_13(ub, 160u);
- var o : mat2x4<f32> = tint_symbol_14(ub, 192u);
- var p : mat3x2<f32> = tint_symbol_15(ub, 224u);
- var q : mat3x3<f32> = tint_symbol_16(ub, 256u);
- var r : mat3x4<f32> = tint_symbol_17(ub, 304u);
- var s : mat4x2<f32> = tint_symbol_18(ub, 352u);
- var t : mat4x3<f32> = tint_symbol_19(ub, 384u);
- var u : mat4x4<f32> = tint_symbol_20(ub, 448u);
- var v : array<vec3<f32>, 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<i32> = tint_symbol_3(&(ub), 16u);
+ var e : vec2<u32> = tint_symbol_4(&(ub), 24u);
+ var f : vec2<f32> = tint_symbol_5(&(ub), 32u);
+ var g : vec3<i32> = tint_symbol_6(&(ub), 48u);
+ var h : vec3<u32> = tint_symbol_7(&(ub), 64u);
+ var i : vec3<f32> = tint_symbol_8(&(ub), 80u);
+ var j : vec4<i32> = tint_symbol_9(&(ub), 96u);
+ var k : vec4<u32> = tint_symbol_10(&(ub), 112u);
+ var l : vec4<f32> = tint_symbol_11(&(ub), 128u);
+ var m : mat2x2<f32> = tint_symbol_12(&(ub), 144u);
+ var n : mat2x3<f32> = tint_symbol_13(&(ub), 160u);
+ var o : mat2x4<f32> = tint_symbol_14(&(ub), 192u);
+ var p : mat3x2<f32> = tint_symbol_15(&(ub), 224u);
+ var q : mat3x3<f32> = tint_symbol_16(&(ub), 256u);
+ var r : mat3x4<f32> = tint_symbol_17(&(ub), 304u);
+ var s : mat4x2<f32> = tint_symbol_18(&(ub), 352u);
+ var t : mat4x3<f32> = tint_symbol_19(&(ub), 384u);
+ var u : mat4x4<f32> = tint_symbol_20(&(ub), 448u);
+ var v : array<vec3<f32>, 2> = tint_symbol_21(&(ub), 512u);
}
)";
@@ -695,78 +695,78 @@
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<uniform, UB, read>, 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<uniform, UB, read>, 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<uniform, UB, read>, 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<i32>
+fn tint_symbol_3(@internal(disable_validation__function_parameter) buffer : ptr<uniform, UB, read>, offset : u32) -> vec2<i32>
@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<u32>
+fn tint_symbol_4(@internal(disable_validation__function_parameter) buffer : ptr<uniform, UB, read>, offset : u32) -> vec2<u32>
@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<f32>
+fn tint_symbol_5(@internal(disable_validation__function_parameter) buffer : ptr<uniform, UB, read>, offset : u32) -> vec2<f32>
@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<i32>
+fn tint_symbol_6(@internal(disable_validation__function_parameter) buffer : ptr<uniform, UB, read>, offset : u32) -> vec3<i32>
@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<u32>
+fn tint_symbol_7(@internal(disable_validation__function_parameter) buffer : ptr<uniform, UB, read>, offset : u32) -> vec3<u32>
@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<f32>
+fn tint_symbol_8(@internal(disable_validation__function_parameter) buffer : ptr<uniform, UB, read>, offset : u32) -> vec3<f32>
@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<i32>
+fn tint_symbol_9(@internal(disable_validation__function_parameter) buffer : ptr<uniform, UB, read>, offset : u32) -> vec4<i32>
@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<u32>
+fn tint_symbol_10(@internal(disable_validation__function_parameter) buffer : ptr<uniform, UB, read>, offset : u32) -> vec4<u32>
@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<f32>
+fn tint_symbol_11(@internal(disable_validation__function_parameter) buffer : ptr<uniform, UB, read>, offset : u32) -> vec4<f32>
-fn tint_symbol_12(@internal(disable_validation__ignore_constructible_function_parameter) buffer : UB, offset : u32) -> mat2x2<f32> {
+fn tint_symbol_12(@internal(disable_validation__function_parameter) buffer : ptr<uniform, UB, read>, offset : u32) -> mat2x2<f32> {
return mat2x2<f32>(tint_symbol_5(buffer, (offset + 0u)), tint_symbol_5(buffer, (offset + 8u)));
}
-fn tint_symbol_13(@internal(disable_validation__ignore_constructible_function_parameter) buffer : UB, offset : u32) -> mat2x3<f32> {
+fn tint_symbol_13(@internal(disable_validation__function_parameter) buffer : ptr<uniform, UB, read>, offset : u32) -> mat2x3<f32> {
return mat2x3<f32>(tint_symbol_8(buffer, (offset + 0u)), tint_symbol_8(buffer, (offset + 16u)));
}
-fn tint_symbol_14(@internal(disable_validation__ignore_constructible_function_parameter) buffer : UB, offset : u32) -> mat2x4<f32> {
+fn tint_symbol_14(@internal(disable_validation__function_parameter) buffer : ptr<uniform, UB, read>, offset : u32) -> mat2x4<f32> {
return mat2x4<f32>(tint_symbol_11(buffer, (offset + 0u)), tint_symbol_11(buffer, (offset + 16u)));
}
-fn tint_symbol_15(@internal(disable_validation__ignore_constructible_function_parameter) buffer : UB, offset : u32) -> mat3x2<f32> {
+fn tint_symbol_15(@internal(disable_validation__function_parameter) buffer : ptr<uniform, UB, read>, offset : u32) -> mat3x2<f32> {
return mat3x2<f32>(tint_symbol_5(buffer, (offset + 0u)), tint_symbol_5(buffer, (offset + 8u)), tint_symbol_5(buffer, (offset + 16u)));
}
-fn tint_symbol_16(@internal(disable_validation__ignore_constructible_function_parameter) buffer : UB, offset : u32) -> mat3x3<f32> {
+fn tint_symbol_16(@internal(disable_validation__function_parameter) buffer : ptr<uniform, UB, read>, offset : u32) -> mat3x3<f32> {
return mat3x3<f32>(tint_symbol_8(buffer, (offset + 0u)), tint_symbol_8(buffer, (offset + 16u)), tint_symbol_8(buffer, (offset + 32u)));
}
-fn tint_symbol_17(@internal(disable_validation__ignore_constructible_function_parameter) buffer : UB, offset : u32) -> mat3x4<f32> {
+fn tint_symbol_17(@internal(disable_validation__function_parameter) buffer : ptr<uniform, UB, read>, offset : u32) -> mat3x4<f32> {
return mat3x4<f32>(tint_symbol_11(buffer, (offset + 0u)), tint_symbol_11(buffer, (offset + 16u)), tint_symbol_11(buffer, (offset + 32u)));
}
-fn tint_symbol_18(@internal(disable_validation__ignore_constructible_function_parameter) buffer : UB, offset : u32) -> mat4x2<f32> {
+fn tint_symbol_18(@internal(disable_validation__function_parameter) buffer : ptr<uniform, UB, read>, offset : u32) -> mat4x2<f32> {
return mat4x2<f32>(tint_symbol_5(buffer, (offset + 0u)), tint_symbol_5(buffer, (offset + 8u)), tint_symbol_5(buffer, (offset + 16u)), tint_symbol_5(buffer, (offset + 24u)));
}
-fn tint_symbol_19(@internal(disable_validation__ignore_constructible_function_parameter) buffer : UB, offset : u32) -> mat4x3<f32> {
+fn tint_symbol_19(@internal(disable_validation__function_parameter) buffer : ptr<uniform, UB, read>, offset : u32) -> mat4x3<f32> {
return mat4x3<f32>(tint_symbol_8(buffer, (offset + 0u)), tint_symbol_8(buffer, (offset + 16u)), tint_symbol_8(buffer, (offset + 32u)), tint_symbol_8(buffer, (offset + 48u)));
}
-fn tint_symbol_20(@internal(disable_validation__ignore_constructible_function_parameter) buffer : UB, offset : u32) -> mat4x4<f32> {
+fn tint_symbol_20(@internal(disable_validation__function_parameter) buffer : ptr<uniform, UB, read>, offset : u32) -> mat4x4<f32> {
return mat4x4<f32>(tint_symbol_11(buffer, (offset + 0u)), tint_symbol_11(buffer, (offset + 16u)), tint_symbol_11(buffer, (offset + 32u)), tint_symbol_11(buffer, (offset + 48u)));
}
-fn tint_symbol_21(@internal(disable_validation__ignore_constructible_function_parameter) buffer : UB, offset : u32) -> array<vec3<f32>, 2u> {
+fn tint_symbol_21(@internal(disable_validation__function_parameter) buffer : ptr<uniform, UB, read>, offset : u32) -> array<vec3<f32>, 2u> {
var arr : array<vec3<f32>, 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 @@
@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<i32> = tint_symbol_3(ub, 16u);
- var e : vec2<u32> = tint_symbol_4(ub, 24u);
- var f : vec2<f32> = tint_symbol_5(ub, 32u);
- var g : vec3<i32> = tint_symbol_6(ub, 48u);
- var h : vec3<u32> = tint_symbol_7(ub, 64u);
- var i : vec3<f32> = tint_symbol_8(ub, 80u);
- var j : vec4<i32> = tint_symbol_9(ub, 96u);
- var k : vec4<u32> = tint_symbol_10(ub, 112u);
- var l : vec4<f32> = tint_symbol_11(ub, 128u);
- var m : mat2x2<f32> = tint_symbol_12(ub, 144u);
- var n : mat2x3<f32> = tint_symbol_13(ub, 160u);
- var o : mat2x4<f32> = tint_symbol_14(ub, 192u);
- var p : mat3x2<f32> = tint_symbol_15(ub, 224u);
- var q : mat3x3<f32> = tint_symbol_16(ub, 256u);
- var r : mat3x4<f32> = tint_symbol_17(ub, 304u);
- var s : mat4x2<f32> = tint_symbol_18(ub, 352u);
- var t : mat4x3<f32> = tint_symbol_19(ub, 384u);
- var u : mat4x4<f32> = tint_symbol_20(ub, 448u);
- var v : array<vec3<f32>, 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<i32> = tint_symbol_3(&(ub), 16u);
+ var e : vec2<u32> = tint_symbol_4(&(ub), 24u);
+ var f : vec2<f32> = tint_symbol_5(&(ub), 32u);
+ var g : vec3<i32> = tint_symbol_6(&(ub), 48u);
+ var h : vec3<u32> = tint_symbol_7(&(ub), 64u);
+ var i : vec3<f32> = tint_symbol_8(&(ub), 80u);
+ var j : vec4<i32> = tint_symbol_9(&(ub), 96u);
+ var k : vec4<u32> = tint_symbol_10(&(ub), 112u);
+ var l : vec4<f32> = tint_symbol_11(&(ub), 128u);
+ var m : mat2x2<f32> = tint_symbol_12(&(ub), 144u);
+ var n : mat2x3<f32> = tint_symbol_13(&(ub), 160u);
+ var o : mat2x4<f32> = tint_symbol_14(&(ub), 192u);
+ var p : mat3x2<f32> = tint_symbol_15(&(ub), 224u);
+ var q : mat3x3<f32> = tint_symbol_16(&(ub), 256u);
+ var r : mat3x4<f32> = tint_symbol_17(&(ub), 304u);
+ var s : mat4x2<f32> = tint_symbol_18(&(ub), 352u);
+ var t : mat4x3<f32> = tint_symbol_19(&(ub), 384u);
+ var u : mat4x4<f32> = tint_symbol_20(&(ub), 448u);
+ var v : array<vec3<f32>, 2> = tint_symbol_21(&(ub), 512u);
}
@group(0) @binding(0) var<uniform> ub : UB;
@@ -918,96 +918,96 @@
@group(0) @binding(0) var<storage, read_write> sb : SB;
@internal(intrinsic_store_storage_i32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : i32)
+fn tint_symbol(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, 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<storage, SB, read_write>, 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<storage, SB, read_write>, offset : u32, value : f32)
@internal(intrinsic_store_storage_vec2_i32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_3(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : vec2<i32>)
+fn tint_symbol_3(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : vec2<i32>)
@internal(intrinsic_store_storage_vec2_u32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_4(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : vec2<u32>)
+fn tint_symbol_4(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : vec2<u32>)
@internal(intrinsic_store_storage_vec2_f32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_5(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : vec2<f32>)
+fn tint_symbol_5(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : vec2<f32>)
@internal(intrinsic_store_storage_vec3_i32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_6(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : vec3<i32>)
+fn tint_symbol_6(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : vec3<i32>)
@internal(intrinsic_store_storage_vec3_u32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_7(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : vec3<u32>)
+fn tint_symbol_7(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : vec3<u32>)
@internal(intrinsic_store_storage_vec3_f32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_8(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : vec3<f32>)
+fn tint_symbol_8(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : vec3<f32>)
@internal(intrinsic_store_storage_vec4_i32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_9(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : vec4<i32>)
+fn tint_symbol_9(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : vec4<i32>)
@internal(intrinsic_store_storage_vec4_u32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_10(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : vec4<u32>)
+fn tint_symbol_10(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : vec4<u32>)
@internal(intrinsic_store_storage_vec4_f32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_11(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : vec4<f32>)
+fn tint_symbol_11(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : vec4<f32>)
-fn tint_symbol_12(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : mat2x2<f32>) {
+fn tint_symbol_12(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : mat2x2<f32>) {
tint_symbol_5(buffer, (offset + 0u), value[0u]);
tint_symbol_5(buffer, (offset + 8u), value[1u]);
}
-fn tint_symbol_13(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : mat2x3<f32>) {
+fn tint_symbol_13(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : mat2x3<f32>) {
tint_symbol_8(buffer, (offset + 0u), value[0u]);
tint_symbol_8(buffer, (offset + 16u), value[1u]);
}
-fn tint_symbol_14(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : mat2x4<f32>) {
+fn tint_symbol_14(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : mat2x4<f32>) {
tint_symbol_11(buffer, (offset + 0u), value[0u]);
tint_symbol_11(buffer, (offset + 16u), value[1u]);
}
-fn tint_symbol_15(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : mat3x2<f32>) {
+fn tint_symbol_15(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : mat3x2<f32>) {
tint_symbol_5(buffer, (offset + 0u), value[0u]);
tint_symbol_5(buffer, (offset + 8u), value[1u]);
tint_symbol_5(buffer, (offset + 16u), value[2u]);
}
-fn tint_symbol_16(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : mat3x3<f32>) {
+fn tint_symbol_16(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : mat3x3<f32>) {
tint_symbol_8(buffer, (offset + 0u), value[0u]);
tint_symbol_8(buffer, (offset + 16u), value[1u]);
tint_symbol_8(buffer, (offset + 32u), value[2u]);
}
-fn tint_symbol_17(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : mat3x4<f32>) {
+fn tint_symbol_17(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : mat3x4<f32>) {
tint_symbol_11(buffer, (offset + 0u), value[0u]);
tint_symbol_11(buffer, (offset + 16u), value[1u]);
tint_symbol_11(buffer, (offset + 32u), value[2u]);
}
-fn tint_symbol_18(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : mat4x2<f32>) {
+fn tint_symbol_18(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : mat4x2<f32>) {
tint_symbol_5(buffer, (offset + 0u), value[0u]);
tint_symbol_5(buffer, (offset + 8u), value[1u]);
tint_symbol_5(buffer, (offset + 16u), value[2u]);
tint_symbol_5(buffer, (offset + 24u), value[3u]);
}
-fn tint_symbol_19(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : mat4x3<f32>) {
+fn tint_symbol_19(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : mat4x3<f32>) {
tint_symbol_8(buffer, (offset + 0u), value[0u]);
tint_symbol_8(buffer, (offset + 16u), value[1u]);
tint_symbol_8(buffer, (offset + 32u), value[2u]);
tint_symbol_8(buffer, (offset + 48u), value[3u]);
}
-fn tint_symbol_20(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : mat4x4<f32>) {
+fn tint_symbol_20(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : mat4x4<f32>) {
tint_symbol_11(buffer, (offset + 0u), value[0u]);
tint_symbol_11(buffer, (offset + 16u), value[1u]);
tint_symbol_11(buffer, (offset + 32u), value[2u]);
tint_symbol_11(buffer, (offset + 48u), value[3u]);
}
-fn tint_symbol_21(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : array<vec3<f32>, 2u>) {
+fn tint_symbol_21(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : array<vec3<f32>, 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 @@
@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<i32>());
- tint_symbol_4(sb, 24u, vec2<u32>());
- tint_symbol_5(sb, 32u, vec2<f32>());
- tint_symbol_6(sb, 48u, vec3<i32>());
- tint_symbol_7(sb, 64u, vec3<u32>());
- tint_symbol_8(sb, 80u, vec3<f32>());
- tint_symbol_9(sb, 96u, vec4<i32>());
- tint_symbol_10(sb, 112u, vec4<u32>());
- tint_symbol_11(sb, 128u, vec4<f32>());
- tint_symbol_12(sb, 144u, mat2x2<f32>());
- tint_symbol_13(sb, 160u, mat2x3<f32>());
- tint_symbol_14(sb, 192u, mat2x4<f32>());
- tint_symbol_15(sb, 224u, mat3x2<f32>());
- tint_symbol_16(sb, 256u, mat3x3<f32>());
- tint_symbol_17(sb, 304u, mat3x4<f32>());
- tint_symbol_18(sb, 352u, mat4x2<f32>());
- tint_symbol_19(sb, 384u, mat4x3<f32>());
- tint_symbol_20(sb, 448u, mat4x4<f32>());
- tint_symbol_21(sb, 512u, array<vec3<f32>, 2>());
+ tint_symbol(&(sb), 0u, i32());
+ tint_symbol_1(&(sb), 4u, u32());
+ tint_symbol_2(&(sb), 8u, f32());
+ tint_symbol_3(&(sb), 16u, vec2<i32>());
+ tint_symbol_4(&(sb), 24u, vec2<u32>());
+ tint_symbol_5(&(sb), 32u, vec2<f32>());
+ tint_symbol_6(&(sb), 48u, vec3<i32>());
+ tint_symbol_7(&(sb), 64u, vec3<u32>());
+ tint_symbol_8(&(sb), 80u, vec3<f32>());
+ tint_symbol_9(&(sb), 96u, vec4<i32>());
+ tint_symbol_10(&(sb), 112u, vec4<u32>());
+ tint_symbol_11(&(sb), 128u, vec4<f32>());
+ tint_symbol_12(&(sb), 144u, mat2x2<f32>());
+ tint_symbol_13(&(sb), 160u, mat2x3<f32>());
+ tint_symbol_14(&(sb), 192u, mat2x4<f32>());
+ tint_symbol_15(&(sb), 224u, mat3x2<f32>());
+ tint_symbol_16(&(sb), 256u, mat3x3<f32>());
+ tint_symbol_17(&(sb), 304u, mat3x4<f32>());
+ tint_symbol_18(&(sb), 352u, mat4x2<f32>());
+ tint_symbol_19(&(sb), 384u, mat4x3<f32>());
+ tint_symbol_20(&(sb), 448u, mat4x4<f32>());
+ tint_symbol_21(&(sb), 512u, array<vec3<f32>, 2>());
}
)";
@@ -1104,96 +1104,96 @@
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<storage, SB, read_write>, 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<storage, SB, read_write>, 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<storage, SB, read_write>, offset : u32, value : f32)
@internal(intrinsic_store_storage_vec2_i32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_3(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : vec2<i32>)
+fn tint_symbol_3(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : vec2<i32>)
@internal(intrinsic_store_storage_vec2_u32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_4(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : vec2<u32>)
+fn tint_symbol_4(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : vec2<u32>)
@internal(intrinsic_store_storage_vec2_f32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_5(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : vec2<f32>)
+fn tint_symbol_5(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : vec2<f32>)
@internal(intrinsic_store_storage_vec3_i32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_6(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : vec3<i32>)
+fn tint_symbol_6(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : vec3<i32>)
@internal(intrinsic_store_storage_vec3_u32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_7(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : vec3<u32>)
+fn tint_symbol_7(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : vec3<u32>)
@internal(intrinsic_store_storage_vec3_f32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_8(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : vec3<f32>)
+fn tint_symbol_8(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : vec3<f32>)
@internal(intrinsic_store_storage_vec4_i32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_9(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : vec4<i32>)
+fn tint_symbol_9(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : vec4<i32>)
@internal(intrinsic_store_storage_vec4_u32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_10(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : vec4<u32>)
+fn tint_symbol_10(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : vec4<u32>)
@internal(intrinsic_store_storage_vec4_f32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_11(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : vec4<f32>)
+fn tint_symbol_11(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : vec4<f32>)
-fn tint_symbol_12(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : mat2x2<f32>) {
+fn tint_symbol_12(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : mat2x2<f32>) {
tint_symbol_5(buffer, (offset + 0u), value[0u]);
tint_symbol_5(buffer, (offset + 8u), value[1u]);
}
-fn tint_symbol_13(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : mat2x3<f32>) {
+fn tint_symbol_13(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : mat2x3<f32>) {
tint_symbol_8(buffer, (offset + 0u), value[0u]);
tint_symbol_8(buffer, (offset + 16u), value[1u]);
}
-fn tint_symbol_14(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : mat2x4<f32>) {
+fn tint_symbol_14(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : mat2x4<f32>) {
tint_symbol_11(buffer, (offset + 0u), value[0u]);
tint_symbol_11(buffer, (offset + 16u), value[1u]);
}
-fn tint_symbol_15(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : mat3x2<f32>) {
+fn tint_symbol_15(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : mat3x2<f32>) {
tint_symbol_5(buffer, (offset + 0u), value[0u]);
tint_symbol_5(buffer, (offset + 8u), value[1u]);
tint_symbol_5(buffer, (offset + 16u), value[2u]);
}
-fn tint_symbol_16(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : mat3x3<f32>) {
+fn tint_symbol_16(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : mat3x3<f32>) {
tint_symbol_8(buffer, (offset + 0u), value[0u]);
tint_symbol_8(buffer, (offset + 16u), value[1u]);
tint_symbol_8(buffer, (offset + 32u), value[2u]);
}
-fn tint_symbol_17(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : mat3x4<f32>) {
+fn tint_symbol_17(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : mat3x4<f32>) {
tint_symbol_11(buffer, (offset + 0u), value[0u]);
tint_symbol_11(buffer, (offset + 16u), value[1u]);
tint_symbol_11(buffer, (offset + 32u), value[2u]);
}
-fn tint_symbol_18(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : mat4x2<f32>) {
+fn tint_symbol_18(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : mat4x2<f32>) {
tint_symbol_5(buffer, (offset + 0u), value[0u]);
tint_symbol_5(buffer, (offset + 8u), value[1u]);
tint_symbol_5(buffer, (offset + 16u), value[2u]);
tint_symbol_5(buffer, (offset + 24u), value[3u]);
}
-fn tint_symbol_19(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : mat4x3<f32>) {
+fn tint_symbol_19(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : mat4x3<f32>) {
tint_symbol_8(buffer, (offset + 0u), value[0u]);
tint_symbol_8(buffer, (offset + 16u), value[1u]);
tint_symbol_8(buffer, (offset + 32u), value[2u]);
tint_symbol_8(buffer, (offset + 48u), value[3u]);
}
-fn tint_symbol_20(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : mat4x4<f32>) {
+fn tint_symbol_20(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : mat4x4<f32>) {
tint_symbol_11(buffer, (offset + 0u), value[0u]);
tint_symbol_11(buffer, (offset + 16u), value[1u]);
tint_symbol_11(buffer, (offset + 32u), value[2u]);
tint_symbol_11(buffer, (offset + 48u), value[3u]);
}
-fn tint_symbol_21(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : array<vec3<f32>, 2u>) {
+fn tint_symbol_21(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : array<vec3<f32>, 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 @@
@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<i32>());
- tint_symbol_4(sb, 24u, vec2<u32>());
- tint_symbol_5(sb, 32u, vec2<f32>());
- tint_symbol_6(sb, 48u, vec3<i32>());
- tint_symbol_7(sb, 64u, vec3<u32>());
- tint_symbol_8(sb, 80u, vec3<f32>());
- tint_symbol_9(sb, 96u, vec4<i32>());
- tint_symbol_10(sb, 112u, vec4<u32>());
- tint_symbol_11(sb, 128u, vec4<f32>());
- tint_symbol_12(sb, 144u, mat2x2<f32>());
- tint_symbol_13(sb, 160u, mat2x3<f32>());
- tint_symbol_14(sb, 192u, mat2x4<f32>());
- tint_symbol_15(sb, 224u, mat3x2<f32>());
- tint_symbol_16(sb, 256u, mat3x3<f32>());
- tint_symbol_17(sb, 304u, mat3x4<f32>());
- tint_symbol_18(sb, 352u, mat4x2<f32>());
- tint_symbol_19(sb, 384u, mat4x3<f32>());
- tint_symbol_20(sb, 448u, mat4x4<f32>());
- tint_symbol_21(sb, 512u, array<vec3<f32>, 2>());
+ tint_symbol(&(sb), 0u, i32());
+ tint_symbol_1(&(sb), 4u, u32());
+ tint_symbol_2(&(sb), 8u, f32());
+ tint_symbol_3(&(sb), 16u, vec2<i32>());
+ tint_symbol_4(&(sb), 24u, vec2<u32>());
+ tint_symbol_5(&(sb), 32u, vec2<f32>());
+ tint_symbol_6(&(sb), 48u, vec3<i32>());
+ tint_symbol_7(&(sb), 64u, vec3<u32>());
+ tint_symbol_8(&(sb), 80u, vec3<f32>());
+ tint_symbol_9(&(sb), 96u, vec4<i32>());
+ tint_symbol_10(&(sb), 112u, vec4<u32>());
+ tint_symbol_11(&(sb), 128u, vec4<f32>());
+ tint_symbol_12(&(sb), 144u, mat2x2<f32>());
+ tint_symbol_13(&(sb), 160u, mat2x3<f32>());
+ tint_symbol_14(&(sb), 192u, mat2x4<f32>());
+ tint_symbol_15(&(sb), 224u, mat3x2<f32>());
+ tint_symbol_16(&(sb), 256u, mat3x3<f32>());
+ tint_symbol_17(&(sb), 304u, mat3x4<f32>());
+ tint_symbol_18(&(sb), 352u, mat4x2<f32>());
+ tint_symbol_19(&(sb), 384u, mat4x3<f32>());
+ tint_symbol_20(&(sb), 448u, mat4x4<f32>());
+ tint_symbol_21(&(sb), 512u, array<vec3<f32>, 2>());
}
@group(0) @binding(0) var<storage, read_write> sb : SB;
@@ -1323,78 +1323,78 @@
@group(0) @binding(0) var<storage, read_write> sb : SB;
@internal(intrinsic_load_storage_i32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> i32
+fn tint_symbol_1(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, 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<storage, SB, read_write>, 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<storage, SB, read_write>, 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<i32>
+fn tint_symbol_4(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> vec2<i32>
@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<u32>
+fn tint_symbol_5(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> vec2<u32>
@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<f32>
+fn tint_symbol_6(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> vec2<f32>
@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<i32>
+fn tint_symbol_7(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> vec3<i32>
@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<u32>
+fn tint_symbol_8(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> vec3<u32>
@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<f32>
+fn tint_symbol_9(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> vec3<f32>
@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<i32>
+fn tint_symbol_10(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> vec4<i32>
@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<u32>
+fn tint_symbol_11(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> vec4<u32>
@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<f32>
+fn tint_symbol_12(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> vec4<f32>
-fn tint_symbol_13(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> mat2x2<f32> {
+fn tint_symbol_13(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> mat2x2<f32> {
return mat2x2<f32>(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<f32> {
+fn tint_symbol_14(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> mat2x3<f32> {
return mat2x3<f32>(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<f32> {
+fn tint_symbol_15(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> mat2x4<f32> {
return mat2x4<f32>(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<f32> {
+fn tint_symbol_16(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> mat3x2<f32> {
return mat3x2<f32>(tint_symbol_6(buffer, (offset + 0u)), tint_symbol_6(buffer, (offset + 8u)), tint_symbol_6(buffer, (offset + 16u)));
}
-fn tint_symbol_17(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> mat3x3<f32> {
+fn tint_symbol_17(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> mat3x3<f32> {
return mat3x3<f32>(tint_symbol_9(buffer, (offset + 0u)), tint_symbol_9(buffer, (offset + 16u)), tint_symbol_9(buffer, (offset + 32u)));
}
-fn tint_symbol_18(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> mat3x4<f32> {
+fn tint_symbol_18(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> mat3x4<f32> {
return mat3x4<f32>(tint_symbol_12(buffer, (offset + 0u)), tint_symbol_12(buffer, (offset + 16u)), tint_symbol_12(buffer, (offset + 32u)));
}
-fn tint_symbol_19(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> mat4x2<f32> {
+fn tint_symbol_19(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> mat4x2<f32> {
return mat4x2<f32>(tint_symbol_6(buffer, (offset + 0u)), tint_symbol_6(buffer, (offset + 8u)), tint_symbol_6(buffer, (offset + 16u)), tint_symbol_6(buffer, (offset + 24u)));
}
-fn tint_symbol_20(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> mat4x3<f32> {
+fn tint_symbol_20(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> mat4x3<f32> {
return mat4x3<f32>(tint_symbol_9(buffer, (offset + 0u)), tint_symbol_9(buffer, (offset + 16u)), tint_symbol_9(buffer, (offset + 32u)), tint_symbol_9(buffer, (offset + 48u)));
}
-fn tint_symbol_21(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> mat4x4<f32> {
+fn tint_symbol_21(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> mat4x4<f32> {
return mat4x4<f32>(tint_symbol_12(buffer, (offset + 0u)), tint_symbol_12(buffer, (offset + 16u)), tint_symbol_12(buffer, (offset + 32u)), tint_symbol_12(buffer, (offset + 48u)));
}
-fn tint_symbol_22(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> array<vec3<f32>, 2u> {
+fn tint_symbol_22(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> array<vec3<f32>, 2u> {
var arr : array<vec3<f32>, 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 @@
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<storage, SB, read_write>, 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 @@
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<storage, SB, read_write>, 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<storage, SB, read_write>, 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<storage, SB, read_write>, 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<i32>
+fn tint_symbol_4(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> vec2<i32>
@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<u32>
+fn tint_symbol_5(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> vec2<u32>
@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<f32>
+fn tint_symbol_6(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> vec2<f32>
@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<i32>
+fn tint_symbol_7(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> vec3<i32>
@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<u32>
+fn tint_symbol_8(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> vec3<u32>
@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<f32>
+fn tint_symbol_9(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> vec3<f32>
@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<i32>
+fn tint_symbol_10(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> vec4<i32>
@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<u32>
+fn tint_symbol_11(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> vec4<u32>
@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<f32>
+fn tint_symbol_12(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> vec4<f32>
-fn tint_symbol_13(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> mat2x2<f32> {
+fn tint_symbol_13(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> mat2x2<f32> {
return mat2x2<f32>(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<f32> {
+fn tint_symbol_14(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> mat2x3<f32> {
return mat2x3<f32>(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<f32> {
+fn tint_symbol_15(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> mat2x4<f32> {
return mat2x4<f32>(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<f32> {
+fn tint_symbol_16(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> mat3x2<f32> {
return mat3x2<f32>(tint_symbol_6(buffer, (offset + 0u)), tint_symbol_6(buffer, (offset + 8u)), tint_symbol_6(buffer, (offset + 16u)));
}
-fn tint_symbol_17(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> mat3x3<f32> {
+fn tint_symbol_17(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> mat3x3<f32> {
return mat3x3<f32>(tint_symbol_9(buffer, (offset + 0u)), tint_symbol_9(buffer, (offset + 16u)), tint_symbol_9(buffer, (offset + 32u)));
}
-fn tint_symbol_18(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> mat3x4<f32> {
+fn tint_symbol_18(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> mat3x4<f32> {
return mat3x4<f32>(tint_symbol_12(buffer, (offset + 0u)), tint_symbol_12(buffer, (offset + 16u)), tint_symbol_12(buffer, (offset + 32u)));
}
-fn tint_symbol_19(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> mat4x2<f32> {
+fn tint_symbol_19(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> mat4x2<f32> {
return mat4x2<f32>(tint_symbol_6(buffer, (offset + 0u)), tint_symbol_6(buffer, (offset + 8u)), tint_symbol_6(buffer, (offset + 16u)), tint_symbol_6(buffer, (offset + 24u)));
}
-fn tint_symbol_20(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> mat4x3<f32> {
+fn tint_symbol_20(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> mat4x3<f32> {
return mat4x3<f32>(tint_symbol_9(buffer, (offset + 0u)), tint_symbol_9(buffer, (offset + 16u)), tint_symbol_9(buffer, (offset + 32u)), tint_symbol_9(buffer, (offset + 48u)));
}
-fn tint_symbol_21(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> mat4x4<f32> {
+fn tint_symbol_21(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> mat4x4<f32> {
return mat4x4<f32>(tint_symbol_12(buffer, (offset + 0u)), tint_symbol_12(buffer, (offset + 16u)), tint_symbol_12(buffer, (offset + 32u)), tint_symbol_12(buffer, (offset + 48u)));
}
-fn tint_symbol_22(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> array<vec3<f32>, 2u> {
+fn tint_symbol_22(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32) -> array<vec3<f32>, 2u> {
var arr : array<vec3<f32>, 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 @@
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<storage, SB, read_write>, 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<storage, read_write> sb : SB;
@@ -1639,103 +1639,103 @@
@group(0) @binding(0) var<storage, read_write> sb : SB;
@internal(intrinsic_store_storage_i32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_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<storage, SB, read_write>, 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<storage, SB, read_write>, 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<storage, SB, read_write>, 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<i32>)
+fn tint_symbol_4(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : vec2<i32>)
@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<u32>)
+fn tint_symbol_5(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : vec2<u32>)
@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<f32>)
+fn tint_symbol_6(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : vec2<f32>)
@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<i32>)
+fn tint_symbol_7(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : vec3<i32>)
@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<u32>)
+fn tint_symbol_8(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : vec3<u32>)
@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<f32>)
+fn tint_symbol_9(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : vec3<f32>)
@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<i32>)
+fn tint_symbol_10(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : vec4<i32>)
@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<u32>)
+fn tint_symbol_11(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : vec4<u32>)
@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<f32>)
+fn tint_symbol_12(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : vec4<f32>)
-fn tint_symbol_13(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : mat2x2<f32>) {
+fn tint_symbol_13(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : mat2x2<f32>) {
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<f32>) {
+fn tint_symbol_14(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : mat2x3<f32>) {
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<f32>) {
+fn tint_symbol_15(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : mat2x4<f32>) {
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<f32>) {
+fn tint_symbol_16(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : mat3x2<f32>) {
tint_symbol_6(buffer, (offset + 0u), value[0u]);
tint_symbol_6(buffer, (offset + 8u), value[1u]);
tint_symbol_6(buffer, (offset + 16u), value[2u]);
}
-fn tint_symbol_17(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : mat3x3<f32>) {
+fn tint_symbol_17(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : mat3x3<f32>) {
tint_symbol_9(buffer, (offset + 0u), value[0u]);
tint_symbol_9(buffer, (offset + 16u), value[1u]);
tint_symbol_9(buffer, (offset + 32u), value[2u]);
}
-fn tint_symbol_18(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : mat3x4<f32>) {
+fn tint_symbol_18(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : mat3x4<f32>) {
tint_symbol_12(buffer, (offset + 0u), value[0u]);
tint_symbol_12(buffer, (offset + 16u), value[1u]);
tint_symbol_12(buffer, (offset + 32u), value[2u]);
}
-fn tint_symbol_19(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : mat4x2<f32>) {
+fn tint_symbol_19(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : mat4x2<f32>) {
tint_symbol_6(buffer, (offset + 0u), value[0u]);
tint_symbol_6(buffer, (offset + 8u), value[1u]);
tint_symbol_6(buffer, (offset + 16u), value[2u]);
tint_symbol_6(buffer, (offset + 24u), value[3u]);
}
-fn tint_symbol_20(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : mat4x3<f32>) {
+fn tint_symbol_20(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : mat4x3<f32>) {
tint_symbol_9(buffer, (offset + 0u), value[0u]);
tint_symbol_9(buffer, (offset + 16u), value[1u]);
tint_symbol_9(buffer, (offset + 32u), value[2u]);
tint_symbol_9(buffer, (offset + 48u), value[3u]);
}
-fn tint_symbol_21(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : mat4x4<f32>) {
+fn tint_symbol_21(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : mat4x4<f32>) {
tint_symbol_12(buffer, (offset + 0u), value[0u]);
tint_symbol_12(buffer, (offset + 16u), value[1u]);
tint_symbol_12(buffer, (offset + 32u), value[2u]);
tint_symbol_12(buffer, (offset + 48u), value[3u]);
}
-fn tint_symbol_22(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : array<vec3<f32>, 2u>) {
+fn tint_symbol_22(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : array<vec3<f32>, 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<storage, SB, read_write>, 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 @@
@compute @workgroup_size(1)
fn main() {
- tint_symbol(sb, 0u, SB());
+ tint_symbol(&(sb), 0u, SB());
}
)";
@@ -1808,103 +1808,103 @@
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<storage, SB, read_write>, 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<storage, SB, read_write>, 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<storage, SB, read_write>, 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<i32>)
+fn tint_symbol_4(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : vec2<i32>)
@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<u32>)
+fn tint_symbol_5(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : vec2<u32>)
@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<f32>)
+fn tint_symbol_6(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : vec2<f32>)
@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<i32>)
+fn tint_symbol_7(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : vec3<i32>)
@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<u32>)
+fn tint_symbol_8(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : vec3<u32>)
@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<f32>)
+fn tint_symbol_9(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : vec3<f32>)
@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<i32>)
+fn tint_symbol_10(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : vec4<i32>)
@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<u32>)
+fn tint_symbol_11(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : vec4<u32>)
@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<f32>)
+fn tint_symbol_12(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : vec4<f32>)
-fn tint_symbol_13(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : mat2x2<f32>) {
+fn tint_symbol_13(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : mat2x2<f32>) {
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<f32>) {
+fn tint_symbol_14(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : mat2x3<f32>) {
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<f32>) {
+fn tint_symbol_15(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : mat2x4<f32>) {
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<f32>) {
+fn tint_symbol_16(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : mat3x2<f32>) {
tint_symbol_6(buffer, (offset + 0u), value[0u]);
tint_symbol_6(buffer, (offset + 8u), value[1u]);
tint_symbol_6(buffer, (offset + 16u), value[2u]);
}
-fn tint_symbol_17(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : mat3x3<f32>) {
+fn tint_symbol_17(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : mat3x3<f32>) {
tint_symbol_9(buffer, (offset + 0u), value[0u]);
tint_symbol_9(buffer, (offset + 16u), value[1u]);
tint_symbol_9(buffer, (offset + 32u), value[2u]);
}
-fn tint_symbol_18(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : mat3x4<f32>) {
+fn tint_symbol_18(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : mat3x4<f32>) {
tint_symbol_12(buffer, (offset + 0u), value[0u]);
tint_symbol_12(buffer, (offset + 16u), value[1u]);
tint_symbol_12(buffer, (offset + 32u), value[2u]);
}
-fn tint_symbol_19(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : mat4x2<f32>) {
+fn tint_symbol_19(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : mat4x2<f32>) {
tint_symbol_6(buffer, (offset + 0u), value[0u]);
tint_symbol_6(buffer, (offset + 8u), value[1u]);
tint_symbol_6(buffer, (offset + 16u), value[2u]);
tint_symbol_6(buffer, (offset + 24u), value[3u]);
}
-fn tint_symbol_20(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : mat4x3<f32>) {
+fn tint_symbol_20(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : mat4x3<f32>) {
tint_symbol_9(buffer, (offset + 0u), value[0u]);
tint_symbol_9(buffer, (offset + 16u), value[1u]);
tint_symbol_9(buffer, (offset + 32u), value[2u]);
tint_symbol_9(buffer, (offset + 48u), value[3u]);
}
-fn tint_symbol_21(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : mat4x4<f32>) {
+fn tint_symbol_21(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : mat4x4<f32>) {
tint_symbol_12(buffer, (offset + 0u), value[0u]);
tint_symbol_12(buffer, (offset + 16u), value[1u]);
tint_symbol_12(buffer, (offset + 32u), value[2u]);
tint_symbol_12(buffer, (offset + 48u), value[3u]);
}
-fn tint_symbol_22(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, value : array<vec3<f32>, 2u>) {
+fn tint_symbol_22(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read_write>, offset : u32, value : array<vec3<f32>, 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<storage, SB, read_write>, 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 @@
@compute @workgroup_size(1)
fn main() {
- tint_symbol(sb, 0u, SB());
+ tint_symbol(&(sb), 0u, SB());
}
@group(0) @binding(0) var<storage, read_write> sb : SB;
@@ -2028,11 +2028,11 @@
@group(0) @binding(0) var<storage, read_write> 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<storage, SB, read_write>, 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 @@
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<storage, SB, read_write>, 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<storage, read_write> sb : SB;
@@ -2164,14 +2164,14 @@
@group(0) @binding(0) var<storage, read_write> 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<storage, SB, read_write>, 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 @@
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<storage, SB, read_write>, 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<storage, read_write> sb : SB;
@@ -2318,14 +2318,14 @@
@group(0) @binding(0) var<storage, read_write> 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<storage, SB, read_write>, 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 @@
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<storage, SB, read_write>, 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<storage, read_write> sb : SB;
@@ -2467,34 +2467,34 @@
@group(0) @binding(0) var<storage, read_write> 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<storage, SB, read_write>, 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<storage, SB, read_write>, 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<storage, SB, read_write>, 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<storage, SB, read_write>, 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<storage, SB, read_write>, 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<storage, SB, read_write>, 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<storage, SB, read_write>, 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<storage, SB, read_write>, 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<storage, SB, read_write>, 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<storage, SB, read_write>, offset : u32, param_1 : i32) -> i32
struct atomic_compare_exchange_weak_ret_type {
old_value : i32,
@@ -2502,37 +2502,37 @@
}
@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<storage, SB, read_write>, 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<storage, SB, read_write>, 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<storage, SB, read_write>, 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<storage, SB, read_write>, 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<storage, SB, read_write>, 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<storage, SB, read_write>, 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<storage, SB, read_write>, 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<storage, SB, read_write>, 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<storage, SB, read_write>, 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<storage, SB, read_write>, 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<storage, SB, read_write>, offset : u32, param_1 : u32) -> u32
struct atomic_compare_exchange_weak_ret_type_1 {
old_value : u32,
@@ -2540,32 +2540,32 @@
}
@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<storage, SB, read_write>, 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 @@
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<storage, SB, read_write>, 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<storage, SB, read_write>, 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<storage, SB, read_write>, 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<storage, SB, read_write>, 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<storage, SB, read_write>, 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<storage, SB, read_write>, 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<storage, SB, read_write>, 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<storage, SB, read_write>, 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<storage, SB, read_write>, 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<storage, SB, read_write>, offset : u32, param_1 : i32) -> i32
struct atomic_compare_exchange_weak_ret_type {
old_value : i32,
@@ -2649,37 +2649,37 @@
}
@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<storage, SB, read_write>, 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<storage, SB, read_write>, 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<storage, SB, read_write>, 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<storage, SB, read_write>, 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<storage, SB, read_write>, 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<storage, SB, read_write>, 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<storage, SB, read_write>, 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<storage, SB, read_write>, 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<storage, SB, read_write>, 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<storage, SB, read_write>, 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<storage, SB, read_write>, offset : u32, param_1 : u32) -> u32
struct atomic_compare_exchange_weak_ret_type_1 {
old_value : u32,
@@ -2687,32 +2687,32 @@
}
@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<storage, SB, read_write>, 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<storage, read_write> sb : SB;
diff --git a/src/tint/transform/manager.cc b/src/tint/transform/manager.cc
index e5f7682..4e83320 100644
--- a/src/tint/transform/manager.cc
+++ b/src/tint/transform/manager.cc
@@ -49,7 +49,8 @@
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 19af4fa..90b50ac 100644
--- a/src/tint/writer/hlsl/generator_impl.cc
+++ b/src/tint/writer/hlsl/generator_impl.cc
@@ -2776,14 +2776,25 @@
first = false;
auto const* type = v->Type();
+ auto storage_class = ast::StorageClass::kNone;
+ auto access = ast::Access::kUndefined;
if (auto* ptr = type->As<sem::Pointer>()) {
- // 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 @@
// 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;
}
