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;
             }