tint: Add DirectVariableAccess transform

Enables the 'chromium_experimental_full_ptr_parameters' extension to
allow passing of uniform, storage and workgroup  address-spaced
pointers as parameters, as well as pointers into sub-objects.

Bug: tint:1758
Change-Id: I8c85e6104ef4f2b9a177dec2857b1bf7f5148212
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/103860
Commit-Queue: Ben Clayton <bclayton@google.com>
Kokoro: Ben Clayton <bclayton@google.com>
Reviewed-by: James Price <jrprice@google.com>

src/tint/BUILD.gn

@@ -502,6 +502,8 @@ libtint_source_set("libtint_core_all_src") {
     "transform/decompose_strided_matrix.h",
     "transform/demote_to_helper.cc",
     "transform/demote_to_helper.h",
+    "transform/direct_variable_access.cc",
+    "transform/direct_variable_access.h",
     "transform/disable_uniformity_analysis.cc",
     "transform/disable_uniformity_analysis.h",
     "transform/expand_compound_assignment.cc",
@@ -1227,6 +1229,7 @@ if (tint_build_unittests) {
       "transform/decompose_strided_array_test.cc",
       "transform/decompose_strided_matrix_test.cc",
       "transform/demote_to_helper_test.cc",
+      "transform/direct_variable_access_test.cc",
       "transform/disable_uniformity_analysis_test.cc",
       "transform/expand_compound_assignment_test.cc",
       "transform/first_index_offset_test.cc",

src/tint/CMakeLists.txt

@@ -425,6 +425,8 @@ list(APPEND TINT_LIB_SRCS
   transform/decompose_strided_matrix.h
   transform/demote_to_helper.cc
   transform/demote_to_helper.h
+  transform/direct_variable_access.cc
+  transform/direct_variable_access.h
   transform/disable_uniformity_analysis.cc
   transform/disable_uniformity_analysis.h
   transform/expand_compound_assignment.cc
@@ -1195,6 +1197,7 @@ if(TINT_BUILD_TESTS)
       transform/decompose_strided_array_test.cc
       transform/decompose_strided_matrix_test.cc
       transform/demote_to_helper_test.cc
+      transform/direct_variable_access_test.cc
       transform/disable_uniformity_analysis_test.cc
       transform/expand_compound_assignment_test.cc
       transform/first_index_offset_test.cc

src/tint/transform/direct_variable_access.h

@@ -0,0 +1,74 @@
+// Copyright 2022 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#ifndef SRC_TINT_TRANSFORM_DIRECT_VARIABLE_ACCESS_H_
+#define SRC_TINT_TRANSFORM_DIRECT_VARIABLE_ACCESS_H_
+
+#include "src/tint/transform/transform.h"
+
+namespace tint::transform {
+
+/// DirectVariableAccess is a transform that allows usage of pointer parameters in the 'storage',
+/// 'uniform' and 'workgroup' address space, and passing of pointers to sub-objects. These pointers
+/// are only allowed by the resolver when the `chromium_experimental_full_ptr_parameters` extension
+/// is enabled.
+///
+/// DirectVariableAccess works by creating specializations of functions that have pointer
+/// parameters, one specialization for each pointer argument's unique access chain 'shape' from a
+/// unique variable. Calls to specialized functions are transformed so that the pointer arguments
+/// are replaced with an array of access-chain indicies, and if the pointer is in the 'function' or
+/// 'private' address space, also with a pointer to the root object. For more information, see the
+/// comments in src/tint/transform/direct_variable_access.cc.
+///
+/// @note DirectVariableAccess requires the transform::Unshadow transform to have been run first.
+class DirectVariableAccess final : public Castable<DirectVariableAccess, Transform> {
+  public:
+    /// Constructor
+    DirectVariableAccess();
+    /// Destructor
+    ~DirectVariableAccess() override;
+
+    /// Options adjusts the behaviour of the transform.
+    struct Options {
+        /// If true, then 'private' sub-object pointer arguments will be transformed.
+        bool transform_private = false;
+        /// If true, then 'function' sub-object pointer arguments will be transformed.
+        bool transform_function = false;
+    };
+
+    /// Config is consumed by the DirectVariableAccess transform.
+    /// Config specifies the behavior of the transform.
+    struct Config final : public Castable<Data, transform::Data> {
+        /// Constructor
+        /// @param options behavior of the transform
+        explicit Config(const Options& options);
+        /// Destructor
+        ~Config() override;
+
+        /// The transform behavior options
+        const Options options;
+    };
+
+    /// @copydoc Transform::Apply
+    ApplyResult Apply(const Program* program,
+                      const DataMap& inputs,
+                      DataMap& outputs) const override;
+
+  private:
+    struct State;
+};
+
+}  // namespace tint::transform
+
+#endif  // SRC_TINT_TRANSFORM_DIRECT_VARIABLE_ACCESS_H_

src/tint/transform/unshadow.h

@@ -19,8 +19,7 @@
 
 namespace tint::transform {
 
-/// Unshadow is a Transform that renames any variables that shadow another
+/// Unshadow is a Transform that renames any variables that shadow another variable.
-/// variable.
 class Unshadow final : public Castable<Unshadow, Transform> {
   public:
     /// Constructor

src/tint/utils/block_allocator.h

@@ -39,6 +39,7 @@ class BlockAllocator {
         static constexpr size_t kMax = 32;
         std::array<T*, kMax> ptrs;
         Pointers* next;
+        Pointers* prev;
     };
 
     /// Block is linked list of memory blocks.
@@ -55,7 +56,7 @@ class BlockAllocator {
     class TView;
 
     /// An iterator for the objects owned by the BlockAllocator.
-    template <bool IS_CONST>
+    template <bool IS_CONST, bool FORWARD>
     class TIterator {
         using PointerTy = std::conditional_t<IS_CONST, const T*, T*>;
 
@@ -72,15 +73,24 @@ class BlockAllocator {
         /// @returns true if this iterator is not equal to other
         bool operator!=(const TIterator& other) const { return !(*this == other); }
 
-        /// Advances the iterator
+        /// Progress the iterator forward one element
         /// @returns this iterator
         TIterator& operator++() {
-            if (ptrs != nullptr) {
+            if (FORWARD) {
-                ++idx;
+                ProgressForward();
-                if (idx == Pointers::kMax) {
+            } else {
-                    idx = 0;
+                ProgressBackwards();
-                    ptrs = ptrs->next;
             }
+            return *this;
+        }
+
+        /// Progress the iterator backwards one element
+        /// @returns this iterator
+        TIterator& operator--() {
+            if (FORWARD) {
+                ProgressBackwards();
+            } else {
+                ProgressForward();
             }
             return *this;
         }
@@ -92,6 +102,27 @@ class BlockAllocator {
         friend TView<IS_CONST>;  // Keep internal iterator impl private.
         explicit TIterator(const Pointers* p, size_t i) : ptrs(p), idx(i) {}
 
+        /// Progresses the iterator forwards
+        void ProgressForward() {
+            if (ptrs != nullptr) {
+                ++idx;
+                if (idx == Pointers::kMax) {
+                    idx = 0;
+                    ptrs = ptrs->next;
+                }
+            }
+        }
+        /// Progresses the iterator backwards
+        void ProgressBackwards() {
+            if (ptrs != nullptr) {
+                if (idx == 0) {
+                    idx = Pointers::kMax - 1;
+                    ptrs = ptrs->prev;
+                }
+                --idx;
+            }
+        }
+
         const Pointers* ptrs;
         size_t idx;
     };
@@ -102,16 +133,25 @@ class BlockAllocator {
     class TView {
       public:
         /// @returns an iterator to the beginning of the view
-        TIterator<IS_CONST> begin() const {
+        TIterator<IS_CONST, true> begin() const {
-            return TIterator<IS_CONST>{allocator_->data.pointers.root, 0};
+            return TIterator<IS_CONST, true>{allocator_->data.pointers.root, 0};
         }
 
         /// @returns an iterator to the end of the view
-        TIterator<IS_CONST> end() const {
+        TIterator<IS_CONST, true> end() const {
             return allocator_->data.pointers.current_index >= Pointers::kMax
-                       ? TIterator<IS_CONST>(nullptr, 0)
+                       ? TIterator<IS_CONST, true>{nullptr, 0}
-                       : TIterator<IS_CONST>(allocator_->data.pointers.current,
+                       : TIterator<IS_CONST, true>{allocator_->data.pointers.current,
-                                             allocator_->data.pointers.current_index);
+                                                   allocator_->data.pointers.current_index};
+        }
+
+        /// @returns an iterator to the beginning of the view
+        TIterator<IS_CONST, false> rbegin() const { return TIterator<IS_CONST, false>{nullptr, 0}; }
+
+        /// @returns an iterator to the end of the view
+        TIterator<IS_CONST, false> rend() const {
+            return TIterator<IS_CONST, false>{allocator_->data.pointers.current,
+                                              allocator_->data.pointers.current_index};
         }
 
       private:
@@ -121,11 +161,17 @@ class BlockAllocator {
     };
 
   public:
-    /// An iterator type over the objects of the BlockAllocator
+    /// A forward-iterator type over the objects of the BlockAllocator
-    using Iterator = TIterator<false>;
+    using Iterator = TIterator</* const */ false, /* forward */ true>;
 
-    /// An immutable iterator type over the objects of the BlockAllocator
+    /// An immutable forward-iterator type over the objects of the BlockAllocator
-    using ConstIterator = TIterator<true>;
+    using ConstIterator = TIterator</* const */ true, /* forward */ true>;
+
+    /// A reverse-iterator type over the objects of the BlockAllocator
+    using ReverseIterator = TIterator</* const */ false, /* forward */ false>;
+
+    /// An immutable reverse-iterator type over the objects of the BlockAllocator
+    using ReverseConstIterator = TIterator</* const */ true, /* forward */ false>;
 
     /// View provides begin() and end() methods for looping over the objects owned by the
     /// BlockAllocator.
@@ -248,6 +294,7 @@ class BlockAllocator {
                 return;  // out of memory
             }
             pointers.current->next = nullptr;
+            pointers.current->prev = prev_pointers;
             pointers.current_index = 0;
 
             if (prev_pointers) {

src/tint/writer/glsl/generator_impl.cc

@@ -54,6 +54,7 @@
 #include "src/tint/transform/combine_samplers.h"
 #include "src/tint/transform/decompose_memory_access.h"
 #include "src/tint/transform/demote_to_helper.h"
+#include "src/tint/transform/direct_variable_access.h"
 #include "src/tint/transform/disable_uniformity_analysis.h"
 #include "src/tint/transform/expand_compound_assignment.h"
 #include "src/tint/transform/manager.h"
@@ -209,7 +210,8 @@ SanitizedResult Sanitize(const Program* in,
     manager.Add<transform::Renamer>();
     data.Add<transform::Renamer::Config>(transform::Renamer::Target::kGlslKeywords,
                                          /* preserve_unicode */ false);
-    manager.Add<transform::Unshadow>();
+    manager.Add<transform::Unshadow>();  // Must come before DirectVariableAccess
+    manager.Add<transform::DirectVariableAccess>();
 
     if (!options.disable_workgroup_init) {
         // ZeroInitWorkgroupMemory must come before CanonicalizeEntryPointIO as

src/tint/writer/hlsl/generator_impl.cc

@@ -54,6 +54,7 @@
 #include "src/tint/transform/canonicalize_entry_point_io.h"
 #include "src/tint/transform/decompose_memory_access.h"
 #include "src/tint/transform/demote_to_helper.h"
+#include "src/tint/transform/direct_variable_access.h"
 #include "src/tint/transform/disable_uniformity_analysis.h"
 #include "src/tint/transform/expand_compound_assignment.h"
 #include "src/tint/transform/localize_struct_array_assignment.h"
@@ -194,7 +195,9 @@ SanitizedResult Sanitize(const Program* in, const Options& options) {
     }
     manager.Add<transform::MultiplanarExternalTexture>();
 
-    manager.Add<transform::Unshadow>();
+    manager.Add<transform::Unshadow>();  // Must come before DirectVariableAccess
+
+    manager.Add<transform::DirectVariableAccess>();
 
     // LocalizeStructArrayAssignment must come after:
     // * SimplifyPointers, because it assumes assignment to arrays in structs are
@@ -291,6 +294,7 @@ bool GeneratorImpl::Generate() {
                                   utils::Vector{
                                       ast::Extension::kChromiumDisableUniformityAnalysis,
                                       ast::Extension::kChromiumExperimentalDp4A,
+                                      ast::Extension::kChromiumExperimentalFullPtrParameters,
                                       ast::Extension::kChromiumExperimentalPushConstant,
                                       ast::Extension::kF16,
                                   })) {

src/tint/writer/msl/generator_impl.cc

@@ -266,6 +266,7 @@ bool GeneratorImpl::Generate() {
     if (!CheckSupportedExtensions("MSL", program_->AST(), diagnostics_,
                                   utils::Vector{
                                       ast::Extension::kChromiumDisableUniformityAnalysis,
+                                      ast::Extension::kChromiumExperimentalFullPtrParameters,
                                       ast::Extension::kChromiumExperimentalPushConstant,
                                       ast::Extension::kF16,
                                   })) {

src/tint/writer/spirv/builder.cc

@@ -260,6 +260,7 @@ bool Builder::Build() {
                                   utils::Vector{
                                       ast::Extension::kChromiumDisableUniformityAnalysis,
                                       ast::Extension::kChromiumExperimentalDp4A,
+                                      ast::Extension::kChromiumExperimentalFullPtrParameters,
                                       ast::Extension::kChromiumExperimentalPushConstant,
                                       ast::Extension::kF16,
                                   })) {

src/tint/writer/spirv/generator_impl.cc

@@ -22,6 +22,7 @@
 #include "src/tint/transform/builtin_polyfill.h"
 #include "src/tint/transform/canonicalize_entry_point_io.h"
 #include "src/tint/transform/demote_to_helper.h"
+#include "src/tint/transform/direct_variable_access.h"
 #include "src/tint/transform/disable_uniformity_analysis.h"
 #include "src/tint/transform/expand_compound_assignment.h"
 #include "src/tint/transform/for_loop_to_loop.h"
@@ -77,12 +78,21 @@ SanitizedResult Sanitize(const Program* in, const Options& options) {
     }
     manager.Add<transform::MultiplanarExternalTexture>();
 
-    manager.Add<transform::Unshadow>();
+    manager.Add<transform::Unshadow>();  // Must come before DirectVariableAccess
     bool disable_workgroup_init_in_sanitizer =
         options.disable_workgroup_init || options.use_zero_initialize_workgroup_memory_extension;
     if (!disable_workgroup_init_in_sanitizer) {
         manager.Add<transform::ZeroInitWorkgroupMemory>();
     }
+
+    {
+        transform::DirectVariableAccess::Options opts;
+        opts.transform_private = true;
+        opts.transform_function = true;
+        data.Add<transform::DirectVariableAccess::Config>(opts);
+        manager.Add<transform::DirectVariableAccess>();
+    }
+
     manager.Add<transform::RemoveUnreachableStatements>();
     manager.Add<transform::PromoteSideEffectsToDecl>();
     manager.Add<transform::SimplifyPointers>();  // Required for arrayLength()

test/tint/ptr_ref/load/param/function/array_in_struct.wgsl

@@ -0,0 +1,15 @@
+enable chromium_experimental_full_ptr_parameters;
+
+struct str {
+  arr : array<i32, 4>,
+};
+
+fn func(pointer : ptr<function, array<i32, 4>>) -> array<i32, 4> {
+  return *pointer;
+}
+
+@compute @workgroup_size(1)
+fn main() {
+  var F : str;
+  let r = func(&F.arr);
+}

test/tint/ptr_ref/load/param/function/array_in_struct.wgsl.expected.dxc.hlsl

@@ -0,0 +1,15 @@
+struct str {
+  int arr[4];
+};
+
+typedef int func_ret[4];
+func_ret func(inout int pointer[4]) {
+  return pointer;
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+  str F = (str)0;
+  const int r[4] = func(F.arr);
+  return;
+}

test/tint/ptr_ref/load/param/function/array_in_struct.wgsl.expected.fxc.hlsl

@@ -0,0 +1,15 @@
+struct str {
+  int arr[4];
+};
+
+typedef int func_ret[4];
+func_ret func(inout int pointer[4]) {
+  return pointer;
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+  str F = (str)0;
+  const int r[4] = func(F.arr);
+  return;
+}

test/tint/ptr_ref/load/param/function/array_in_struct.wgsl.expected.glsl

@@ -1,19 +1,16 @@
 #version 310 es
 
-void func(int value, inout int pointer) {
+struct str {
-  pointer = value;
+  int arr[4];
-  return;
+};
-}
 
-void main_1() {
+int[4] func(inout int pointer[4]) {
-  int i = 0;
+  return pointer;
-  i = 123;
-  func(123, i);
-  return;
 }
 
 void tint_symbol() {
-  main_1();
+  str F = str(int[4](0, 0, 0, 0));
+  int r[4] = func(F.arr);
 }
 
 layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;

test/tint/ptr_ref/load/param/function/array_in_struct.wgsl.expected.msl

@@ -0,0 +1,30 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+template<typename T, size_t N>
+struct tint_array {
+    const constant T& operator[](size_t i) const constant { return elements[i]; }
+    device T& operator[](size_t i) device { return elements[i]; }
+    const device T& operator[](size_t i) const device { return elements[i]; }
+    thread T& operator[](size_t i) thread { return elements[i]; }
+    const thread T& operator[](size_t i) const thread { return elements[i]; }
+    threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+    const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+    T elements[N];
+};
+
+struct str {
+  tint_array<int, 4> arr;
+};
+
+tint_array<int, 4> func(thread tint_array<int, 4>* const pointer) {
+  return *(pointer);
+}
+
+kernel void tint_symbol() {
+  str F = {};
+  tint_array<int, 4> const r = func(&(F.arr));
+  return;
+}
+

test/tint/ptr_ref/load/param/function/array_in_struct.wgsl.expected.spvasm

@@ -0,0 +1,42 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 24
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %str "str"
+               OpMemberName %str 0 "arr"
+               OpName %func_F_arr "func_F_arr"
+               OpName %pointer "pointer"
+               OpName %main "main"
+               OpName %F "F"
+               OpDecorate %_arr_int_uint_4 ArrayStride 4
+               OpMemberDecorate %str 0 Offset 0
+        %int = OpTypeInt 32 1
+       %uint = OpTypeInt 32 0
+     %uint_4 = OpConstant %uint 4
+%_arr_int_uint_4 = OpTypeArray %int %uint_4
+        %str = OpTypeStruct %_arr_int_uint_4
+%_ptr_Function_str = OpTypePointer Function %str
+          %1 = OpTypeFunction %_arr_int_uint_4 %_ptr_Function_str
+     %uint_0 = OpConstant %uint 0
+%_ptr_Function__arr_int_uint_4 = OpTypePointer Function %_arr_int_uint_4
+       %void = OpTypeVoid
+         %16 = OpTypeFunction %void
+         %21 = OpConstantNull %str
+ %func_F_arr = OpFunction %_arr_int_uint_4 None %1
+    %pointer = OpFunctionParameter %_ptr_Function_str
+         %10 = OpLabel
+         %14 = OpAccessChain %_ptr_Function__arr_int_uint_4 %pointer %uint_0
+         %15 = OpLoad %_arr_int_uint_4 %14
+               OpReturnValue %15
+               OpFunctionEnd
+       %main = OpFunction %void None %16
+         %19 = OpLabel
+          %F = OpVariable %_ptr_Function_str Function %21
+         %22 = OpFunctionCall %_arr_int_uint_4 %func_F_arr %F
+               OpReturn
+               OpFunctionEnd

test/tint/ptr_ref/load/param/function/array_in_struct.wgsl.expected.wgsl

@@ -0,0 +1,15 @@
+enable chromium_experimental_full_ptr_parameters;
+
+struct str {
+  arr : array<i32, 4>,
+}
+
+fn func(pointer : ptr<function, array<i32, 4>>) -> array<i32, 4> {
+  return *(pointer);
+}
+
+@compute @workgroup_size(1)
+fn main() {
+  var F : str;
+  let r = func(&(F.arr));
+}

test/tint/ptr_ref/load/param/function/i32.wgsl

@@ -0,0 +1,9 @@
+fn func(pointer : ptr<function, i32>) -> i32 {
+  return *pointer;
+}
+
+@compute @workgroup_size(1)
+fn main() {
+  var F : i32;
+  let r = func(&F);
+}

test/tint/ptr_ref/load/param/function/i32.wgsl.expected.dxc.hlsl

@@ -0,0 +1,10 @@
+int func(inout int pointer) {
+  return pointer;
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+  int F = 0;
+  const int r = func(F);
+  return;
+}

test/tint/ptr_ref/load/param/function/i32.wgsl.expected.fxc.hlsl

@@ -0,0 +1,10 @@
+int func(inout int pointer) {
+  return pointer;
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+  int F = 0;
+  const int r = func(F);
+  return;
+}

test/tint/ptr_ref/load/param/function/i32.wgsl.expected.glsl

@@ -0,0 +1,16 @@
+#version 310 es
+
+int func(inout int pointer) {
+  return pointer;
+}
+
+void tint_symbol() {
+  int F = 0;
+  int r = func(F);
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  tint_symbol();
+  return;
+}

test/tint/ptr_ref/load/param/function/i32.wgsl.expected.msl

@@ -0,0 +1,13 @@
+#include <metal_stdlib>
+
+using namespace metal;
+int func(thread int* const pointer) {
+  return *(pointer);
+}
+
+kernel void tint_symbol() {
+  int F = 0;
+  int const r = func(&(F));
+  return;
+}
+

test/tint/ptr_ref/load/param/function/i32.wgsl.expected.spvasm

@@ -0,0 +1,31 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 17
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %func "func"
+               OpName %pointer "pointer"
+               OpName %main "main"
+               OpName %F "F"
+        %int = OpTypeInt 32 1
+%_ptr_Function_int = OpTypePointer Function %int
+          %1 = OpTypeFunction %int %_ptr_Function_int
+       %void = OpTypeVoid
+          %9 = OpTypeFunction %void
+         %14 = OpConstantNull %int
+       %func = OpFunction %int None %1
+    %pointer = OpFunctionParameter %_ptr_Function_int
+          %6 = OpLabel
+          %8 = OpLoad %int %pointer
+               OpReturnValue %8
+               OpFunctionEnd
+       %main = OpFunction %void None %9
+         %12 = OpLabel
+          %F = OpVariable %_ptr_Function_int Function %14
+         %15 = OpFunctionCall %int %func %F
+               OpReturn
+               OpFunctionEnd

test/tint/ptr_ref/load/param/function/i32.wgsl.expected.wgsl

@@ -0,0 +1,9 @@
+fn func(pointer : ptr<function, i32>) -> i32 {
+  return *(pointer);
+}
+
+@compute @workgroup_size(1)
+fn main() {
+  var F : i32;
+  let r = func(&(F));
+}

test/tint/ptr_ref/load/param/function/i32_in_struct.wgsl

@@ -0,0 +1,15 @@
+enable chromium_experimental_full_ptr_parameters;
+
+struct str {
+  i : i32,
+};
+
+fn func(pointer : ptr<function, i32>) -> i32 {
+  return *pointer;
+}
+
+@compute @workgroup_size(1)
+fn main() {
+  var F : str;
+  let r = func(&F.i);
+}

test/tint/ptr_ref/load/param/function/i32_in_struct.wgsl.expected.dxc.hlsl

@@ -0,0 +1,14 @@
+struct str {
+  int i;
+};
+
+int func(inout int pointer) {
+  return pointer;
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+  str F = (str)0;
+  const int r = func(F.i);
+  return;
+}

test/tint/ptr_ref/load/param/function/i32_in_struct.wgsl.expected.fxc.hlsl

@@ -0,0 +1,14 @@
+struct str {
+  int i;
+};
+
+int func(inout int pointer) {
+  return pointer;
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+  str F = (str)0;
+  const int r = func(F.i);
+  return;
+}

test/tint/ptr_ref/load/param/function/i32_in_struct.wgsl.expected.glsl

@@ -0,0 +1,20 @@
+#version 310 es
+
+struct str {
+  int i;
+};
+
+int func(inout int pointer) {
+  return pointer;
+}
+
+void tint_symbol() {
+  str F = str(0);
+  int r = func(F.i);
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  tint_symbol();
+  return;
+}

test/tint/ptr_ref/load/param/function/i32_in_struct.wgsl.expected.msl

@@ -0,0 +1,17 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct str {
+  int i;
+};
+
+int func(thread int* const pointer) {
+  return *(pointer);
+}
+
+kernel void tint_symbol() {
+  str F = {};
+  int const r = func(&(F.i));
+  return;
+}
+

test/tint/ptr_ref/load/param/function/i32_in_struct.wgsl.expected.spvasm

@@ -0,0 +1,39 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 22
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %str "str"
+               OpMemberName %str 0 "i"
+               OpName %func_F_i "func_F_i"
+               OpName %pointer "pointer"
+               OpName %main "main"
+               OpName %F "F"
+               OpMemberDecorate %str 0 Offset 0
+        %int = OpTypeInt 32 1
+        %str = OpTypeStruct %int
+%_ptr_Function_str = OpTypePointer Function %str
+          %1 = OpTypeFunction %int %_ptr_Function_str
+       %uint = OpTypeInt 32 0
+     %uint_0 = OpConstant %uint 0
+%_ptr_Function_int = OpTypePointer Function %int
+       %void = OpTypeVoid
+         %14 = OpTypeFunction %void
+         %19 = OpConstantNull %str
+   %func_F_i = OpFunction %int None %1
+    %pointer = OpFunctionParameter %_ptr_Function_str
+          %7 = OpLabel
+         %12 = OpAccessChain %_ptr_Function_int %pointer %uint_0
+         %13 = OpLoad %int %12
+               OpReturnValue %13
+               OpFunctionEnd
+       %main = OpFunction %void None %14
+         %17 = OpLabel
+          %F = OpVariable %_ptr_Function_str Function %19
+         %20 = OpFunctionCall %int %func_F_i %F
+               OpReturn
+               OpFunctionEnd

test/tint/ptr_ref/load/param/function/i32_in_struct.wgsl.expected.wgsl

@@ -0,0 +1,15 @@
+enable chromium_experimental_full_ptr_parameters;
+
+struct str {
+  i : i32,
+}
+
+fn func(pointer : ptr<function, i32>) -> i32 {
+  return *(pointer);
+}
+
+@compute @workgroup_size(1)
+fn main() {
+  var F : str;
+  let r = func(&(F.i));
+}

test/tint/ptr_ref/load/param/function/struct_in_array.wgsl

@@ -0,0 +1,15 @@
+enable chromium_experimental_full_ptr_parameters;
+
+struct str {
+  i : i32,
+};
+
+fn func(pointer : ptr<function, str>) -> str {
+  return *pointer;
+}
+
+@compute @workgroup_size(1)
+fn main() {
+  var F : array<str, 4>;
+  let r = func(&F[2]);
+}

test/tint/ptr_ref/load/param/function/struct_in_array.wgsl.expected.dxc.hlsl

@@ -0,0 +1,14 @@
+struct str {
+  int i;
+};
+
+str func(inout str pointer) {
+  return pointer;
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+  str F[4] = (str[4])0;
+  const str r = func(F[2]);
+  return;
+}

test/tint/ptr_ref/load/param/function/struct_in_array.wgsl.expected.fxc.hlsl

@@ -0,0 +1,14 @@
+struct str {
+  int i;
+};
+
+str func(inout str pointer) {
+  return pointer;
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+  str F[4] = (str[4])0;
+  const str r = func(F[2]);
+  return;
+}

test/tint/ptr_ref/load/param/function/struct_in_array.wgsl.expected.glsl

@@ -0,0 +1,20 @@
+#version 310 es
+
+struct str {
+  int i;
+};
+
+str func(inout str pointer) {
+  return pointer;
+}
+
+void tint_symbol() {
+  str F[4] = str[4](str(0), str(0), str(0), str(0));
+  str r = func(F[2]);
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  tint_symbol();
+  return;
+}

test/tint/ptr_ref/load/param/function/struct_in_array.wgsl.expected.msl

@@ -0,0 +1,30 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+template<typename T, size_t N>
+struct tint_array {
+    const constant T& operator[](size_t i) const constant { return elements[i]; }
+    device T& operator[](size_t i) device { return elements[i]; }
+    const device T& operator[](size_t i) const device { return elements[i]; }
+    thread T& operator[](size_t i) thread { return elements[i]; }
+    const thread T& operator[](size_t i) const thread { return elements[i]; }
+    threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+    const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+    T elements[N];
+};
+
+struct str {
+  int i;
+};
+
+str func(thread str* const pointer) {
+  return *(pointer);
+}
+
+kernel void tint_symbol() {
+  tint_array<str, 4> F = {};
+  str const r = func(&(F[2]));
+  return;
+}
+

test/tint/ptr_ref/load/param/function/struct_in_array.wgsl.expected.spvasm

@@ -0,0 +1,50 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 30
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %str "str"
+               OpMemberName %str 0 "i"
+               OpName %func_F_X "func_F_X"
+               OpName %pointer_base "pointer_base"
+               OpName %pointer_indices "pointer_indices"
+               OpName %main "main"
+               OpName %F "F"
+               OpMemberDecorate %str 0 Offset 0
+               OpDecorate %_arr_str_uint_4 ArrayStride 4
+               OpDecorate %_arr_uint_uint_1 ArrayStride 4
+        %int = OpTypeInt 32 1
+        %str = OpTypeStruct %int
+       %uint = OpTypeInt 32 0
+     %uint_4 = OpConstant %uint 4
+%_arr_str_uint_4 = OpTypeArray %str %uint_4
+%_ptr_Function__arr_str_uint_4 = OpTypePointer Function %_arr_str_uint_4
+     %uint_1 = OpConstant %uint 1
+%_arr_uint_uint_1 = OpTypeArray %uint %uint_1
+          %1 = OpTypeFunction %str %_ptr_Function__arr_str_uint_4 %_arr_uint_uint_1
+         %15 = OpConstantNull %int
+%_ptr_Function_str = OpTypePointer Function %str
+       %void = OpTypeVoid
+         %20 = OpTypeFunction %void
+         %25 = OpConstantNull %_arr_str_uint_4
+     %uint_2 = OpConstant %uint 2
+         %29 = OpConstantComposite %_arr_uint_uint_1 %uint_2
+   %func_F_X = OpFunction %str None %1
+%pointer_base = OpFunctionParameter %_ptr_Function__arr_str_uint_4
+%pointer_indices = OpFunctionParameter %_arr_uint_uint_1
+         %13 = OpLabel
+         %16 = OpCompositeExtract %uint %pointer_indices 0
+         %18 = OpAccessChain %_ptr_Function_str %pointer_base %16
+         %19 = OpLoad %str %18
+               OpReturnValue %19
+               OpFunctionEnd
+       %main = OpFunction %void None %20
+         %23 = OpLabel
+          %F = OpVariable %_ptr_Function__arr_str_uint_4 Function %25
+         %26 = OpFunctionCall %str %func_F_X %F %29
+               OpReturn
+               OpFunctionEnd

test/tint/ptr_ref/load/param/function/struct_in_array.wgsl.expected.wgsl

@@ -0,0 +1,15 @@
+enable chromium_experimental_full_ptr_parameters;
+
+struct str {
+  i : i32,
+}
+
+fn func(pointer : ptr<function, str>) -> str {
+  return *(pointer);
+}
+
+@compute @workgroup_size(1)
+fn main() {
+  var F : array<str, 4>;
+  let r = func(&(F[2]));
+}

test/tint/ptr_ref/load/param/function/vec2_f32_in_mat2x2.wgsl

@@ -0,0 +1,11 @@
+enable chromium_experimental_full_ptr_parameters;
+
+fn func(pointer : ptr<function, vec2<f32>>) -> vec2<f32> {
+  return *pointer;
+}
+
+@compute @workgroup_size(1)
+fn main() {
+  var F : mat2x2<f32>;
+  let r = func(&F[1]);
+}

test/tint/ptr_ref/load/param/function/vec2_f32_in_mat2x2.wgsl.expected.dxc.hlsl

@@ -0,0 +1,10 @@
+float2 func(inout float2 pointer) {
+  return pointer;
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+  float2x2 F = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+  const float2 r = func(F[1]);
+  return;
+}

test/tint/ptr_ref/load/param/function/vec2_f32_in_mat2x2.wgsl.expected.fxc.hlsl

@@ -0,0 +1,10 @@
+float2 func(inout float2 pointer) {
+  return pointer;
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+  float2x2 F = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+  const float2 r = func(F[1]);
+  return;
+}

test/tint/ptr_ref/load/param/function/vec2_f32_in_mat2x2.wgsl.expected.glsl

@@ -0,0 +1,16 @@
+#version 310 es
+
+vec2 func(inout vec2 pointer) {
+  return pointer;
+}
+
+void tint_symbol() {
+  mat2 F = mat2(0.0f, 0.0f, 0.0f, 0.0f);
+  vec2 r = func(F[1]);
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  tint_symbol();
+  return;
+}

test/tint/ptr_ref/load/param/function/vec2_f32_in_mat2x2.wgsl.expected.msl

@@ -0,0 +1,13 @@
+#include <metal_stdlib>
+
+using namespace metal;
+float2 func(thread float2* const pointer) {
+  return *(pointer);
+}
+
+kernel void tint_symbol() {
+  float2x2 F = float2x2(0.0f);
+  float2 const r = func(&(F[1]));
+  return;
+}
+

test/tint/ptr_ref/load/param/function/vec2_f32_in_mat2x2.wgsl.expected.spvasm

@@ -0,0 +1,45 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 29
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %func_F_X "func_F_X"
+               OpName %pointer_base "pointer_base"
+               OpName %pointer_indices "pointer_indices"
+               OpName %main "main"
+               OpName %F "F"
+               OpDecorate %_arr_uint_uint_1 ArrayStride 4
+      %float = OpTypeFloat 32
+    %v2float = OpTypeVector %float 2
+%mat2v2float = OpTypeMatrix %v2float 2
+%_ptr_Function_mat2v2float = OpTypePointer Function %mat2v2float
+       %uint = OpTypeInt 32 0
+     %uint_1 = OpConstant %uint 1
+%_arr_uint_uint_1 = OpTypeArray %uint %uint_1
+          %1 = OpTypeFunction %v2float %_ptr_Function_mat2v2float %_arr_uint_uint_1
+        %int = OpTypeInt 32 1
+         %15 = OpConstantNull %int
+%_ptr_Function_v2float = OpTypePointer Function %v2float
+       %void = OpTypeVoid
+         %20 = OpTypeFunction %void
+         %25 = OpConstantNull %mat2v2float
+         %28 = OpConstantComposite %_arr_uint_uint_1 %uint_1
+   %func_F_X = OpFunction %v2float None %1
+%pointer_base = OpFunctionParameter %_ptr_Function_mat2v2float
+%pointer_indices = OpFunctionParameter %_arr_uint_uint_1
+         %12 = OpLabel
+         %16 = OpCompositeExtract %uint %pointer_indices 0
+         %18 = OpAccessChain %_ptr_Function_v2float %pointer_base %16
+         %19 = OpLoad %v2float %18
+               OpReturnValue %19
+               OpFunctionEnd
+       %main = OpFunction %void None %20
+         %23 = OpLabel
+          %F = OpVariable %_ptr_Function_mat2v2float Function %25
+         %26 = OpFunctionCall %v2float %func_F_X %F %28
+               OpReturn
+               OpFunctionEnd

test/tint/ptr_ref/load/param/function/vec2_f32_in_mat2x2.wgsl.expected.wgsl

@@ -0,0 +1,11 @@
+enable chromium_experimental_full_ptr_parameters;
+
+fn func(pointer : ptr<function, vec2<f32>>) -> vec2<f32> {
+  return *(pointer);
+}
+
+@compute @workgroup_size(1)
+fn main() {
+  var F : mat2x2<f32>;
+  let r = func(&(F[1]));
+}

test/tint/ptr_ref/load/param/function/vec4_f32.wgsl

@@ -0,0 +1,9 @@
+fn func(pointer : ptr<function, vec4<f32>>) -> vec4<f32> {
+  return *pointer;
+}
+
+@compute @workgroup_size(1)
+fn main() {
+  var F : vec4<f32>;
+  let r = func(&F);
+}

test/tint/ptr_ref/load/param/function/vec4_f32.wgsl.expected.dxc.hlsl

@@ -0,0 +1,10 @@
+float4 func(inout float4 pointer) {
+  return pointer;
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+  float4 F = float4(0.0f, 0.0f, 0.0f, 0.0f);
+  const float4 r = func(F);
+  return;
+}

test/tint/ptr_ref/load/param/function/vec4_f32.wgsl.expected.fxc.hlsl

@@ -0,0 +1,10 @@
+float4 func(inout float4 pointer) {
+  return pointer;
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+  float4 F = float4(0.0f, 0.0f, 0.0f, 0.0f);
+  const float4 r = func(F);
+  return;
+}

test/tint/ptr_ref/load/param/function/vec4_f32.wgsl.expected.glsl

@@ -0,0 +1,16 @@
+#version 310 es
+
+vec4 func(inout vec4 pointer) {
+  return pointer;
+}
+
+void tint_symbol() {
+  vec4 F = vec4(0.0f, 0.0f, 0.0f, 0.0f);
+  vec4 r = func(F);
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  tint_symbol();
+  return;
+}

test/tint/ptr_ref/load/param/function/vec4_f32.wgsl.expected.msl

@@ -0,0 +1,13 @@
+#include <metal_stdlib>
+
+using namespace metal;
+float4 func(thread float4* const pointer) {
+  return *(pointer);
+}
+
+kernel void tint_symbol() {
+  float4 F = 0.0f;
+  float4 const r = func(&(F));
+  return;
+}
+

test/tint/ptr_ref/load/param/function/vec4_f32.wgsl.expected.spvasm

@@ -0,0 +1,32 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 18
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %func "func"
+               OpName %pointer "pointer"
+               OpName %main "main"
+               OpName %F "F"
+      %float = OpTypeFloat 32
+    %v4float = OpTypeVector %float 4
+%_ptr_Function_v4float = OpTypePointer Function %v4float
+          %1 = OpTypeFunction %v4float %_ptr_Function_v4float
+       %void = OpTypeVoid
+         %10 = OpTypeFunction %void
+         %15 = OpConstantNull %v4float
+       %func = OpFunction %v4float None %1
+    %pointer = OpFunctionParameter %_ptr_Function_v4float
+          %7 = OpLabel
+          %9 = OpLoad %v4float %pointer
+               OpReturnValue %9
+               OpFunctionEnd
+       %main = OpFunction %void None %10
+         %13 = OpLabel
+          %F = OpVariable %_ptr_Function_v4float Function %15
+         %16 = OpFunctionCall %v4float %func %F
+               OpReturn
+               OpFunctionEnd

test/tint/ptr_ref/load/param/function/vec4_f32.wgsl.expected.wgsl

@@ -0,0 +1,9 @@
+fn func(pointer : ptr<function, vec4<f32>>) -> vec4<f32> {
+  return *(pointer);
+}
+
+@compute @workgroup_size(1)
+fn main() {
+  var F : vec4<f32>;
+  let r = func(&(F));
+}

test/tint/ptr_ref/load/param/function/vec4_f32_in_mat2x4.wgsl

@@ -0,0 +1,11 @@
+enable chromium_experimental_full_ptr_parameters;
+
+fn func(pointer : ptr<function, vec4<f32>>) -> vec4<f32> {
+  return *pointer;
+}
+
+@compute @workgroup_size(1)
+fn main() {
+  var F : mat2x4<f32>;
+  let r = func(&F[1]);
+}

test/tint/ptr_ref/load/param/function/vec4_f32_in_mat2x4.wgsl.expected.dxc.hlsl

@@ -0,0 +1,10 @@
+float4 func(inout float4 pointer) {
+  return pointer;
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+  float2x4 F = float2x4(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f);
+  const float4 r = func(F[1]);
+  return;
+}

test/tint/ptr_ref/load/param/function/vec4_f32_in_mat2x4.wgsl.expected.fxc.hlsl

@@ -0,0 +1,10 @@
+float4 func(inout float4 pointer) {
+  return pointer;
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+  float2x4 F = float2x4(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f);
+  const float4 r = func(F[1]);
+  return;
+}

test/tint/ptr_ref/load/param/function/vec4_f32_in_mat2x4.wgsl.expected.glsl

@@ -0,0 +1,16 @@
+#version 310 es
+
+vec4 func(inout vec4 pointer) {
+  return pointer;
+}
+
+void tint_symbol() {
+  mat2x4 F = mat2x4(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f);
+  vec4 r = func(F[1]);
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  tint_symbol();
+  return;
+}

test/tint/ptr_ref/load/param/function/vec4_f32_in_mat2x4.wgsl.expected.msl

@@ -0,0 +1,13 @@
+#include <metal_stdlib>
+
+using namespace metal;
+float4 func(thread float4* const pointer) {
+  return *(pointer);
+}
+
+kernel void tint_symbol() {
+  float2x4 F = float2x4(0.0f);
+  float4 const r = func(&(F[1]));
+  return;
+}
+

test/tint/ptr_ref/load/param/function/vec4_f32_in_mat2x4.wgsl.expected.spvasm

@@ -0,0 +1,45 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 29
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %func_F_X "func_F_X"
+               OpName %pointer_base "pointer_base"
+               OpName %pointer_indices "pointer_indices"
+               OpName %main "main"
+               OpName %F "F"
+               OpDecorate %_arr_uint_uint_1 ArrayStride 4
+      %float = OpTypeFloat 32
+    %v4float = OpTypeVector %float 4
+%mat2v4float = OpTypeMatrix %v4float 2
+%_ptr_Function_mat2v4float = OpTypePointer Function %mat2v4float
+       %uint = OpTypeInt 32 0
+     %uint_1 = OpConstant %uint 1
+%_arr_uint_uint_1 = OpTypeArray %uint %uint_1
+          %1 = OpTypeFunction %v4float %_ptr_Function_mat2v4float %_arr_uint_uint_1
+        %int = OpTypeInt 32 1
+         %15 = OpConstantNull %int
+%_ptr_Function_v4float = OpTypePointer Function %v4float
+       %void = OpTypeVoid
+         %20 = OpTypeFunction %void
+         %25 = OpConstantNull %mat2v4float
+         %28 = OpConstantComposite %_arr_uint_uint_1 %uint_1
+   %func_F_X = OpFunction %v4float None %1
+%pointer_base = OpFunctionParameter %_ptr_Function_mat2v4float
+%pointer_indices = OpFunctionParameter %_arr_uint_uint_1
+         %12 = OpLabel
+         %16 = OpCompositeExtract %uint %pointer_indices 0
+         %18 = OpAccessChain %_ptr_Function_v4float %pointer_base %16
+         %19 = OpLoad %v4float %18
+               OpReturnValue %19
+               OpFunctionEnd
+       %main = OpFunction %void None %20
+         %23 = OpLabel
+          %F = OpVariable %_ptr_Function_mat2v4float Function %25
+         %26 = OpFunctionCall %v4float %func_F_X %F %28
+               OpReturn
+               OpFunctionEnd

test/tint/ptr_ref/load/param/function/vec4_f32_in_mat2x4.wgsl.expected.wgsl

@@ -0,0 +1,11 @@
+enable chromium_experimental_full_ptr_parameters;
+
+fn func(pointer : ptr<function, vec4<f32>>) -> vec4<f32> {
+  return *(pointer);
+}
+
+@compute @workgroup_size(1)
+fn main() {
+  var F : mat2x4<f32>;
+  let r = func(&(F[1]));
+}

test/tint/ptr_ref/load/param/function/vec4_f32_in_struct.wgsl

@@ -0,0 +1,15 @@
+enable chromium_experimental_full_ptr_parameters;
+
+struct str {
+  i : vec4<f32>,
+};
+
+fn func(pointer : ptr<function, vec4<f32>>) -> vec4<f32> {
+  return *pointer;
+}
+
+@compute @workgroup_size(1)
+fn main() {
+  var F : str;
+  let r = func(&F.i);
+}

test/tint/ptr_ref/load/param/function/vec4_f32_in_struct.wgsl.expected.dxc.hlsl

@@ -0,0 +1,14 @@
+struct str {
+  float4 i;
+};
+
+float4 func(inout float4 pointer) {
+  return pointer;
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+  str F = (str)0;
+  const float4 r = func(F.i);
+  return;
+}

test/tint/ptr_ref/load/param/function/vec4_f32_in_struct.wgsl.expected.fxc.hlsl

@@ -0,0 +1,14 @@
+struct str {
+  float4 i;
+};
+
+float4 func(inout float4 pointer) {
+  return pointer;
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+  str F = (str)0;
+  const float4 r = func(F.i);
+  return;
+}

test/tint/ptr_ref/load/param/function/vec4_f32_in_struct.wgsl.expected.glsl

@@ -0,0 +1,20 @@
+#version 310 es
+
+struct str {
+  vec4 i;
+};
+
+vec4 func(inout vec4 pointer) {
+  return pointer;
+}
+
+void tint_symbol() {
+  str F = str(vec4(0.0f, 0.0f, 0.0f, 0.0f));
+  vec4 r = func(F.i);
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  tint_symbol();
+  return;
+}

test/tint/ptr_ref/load/param/function/vec4_f32_in_struct.wgsl.expected.msl

@@ -0,0 +1,17 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct str {
+  float4 i;
+};
+
+float4 func(thread float4* const pointer) {
+  return *(pointer);
+}
+
+kernel void tint_symbol() {
+  str F = {};
+  float4 const r = func(&(F.i));
+  return;
+}
+

test/tint/ptr_ref/load/param/function/vec4_f32_in_struct.wgsl.expected.spvasm

@@ -0,0 +1,40 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 23
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %str "str"
+               OpMemberName %str 0 "i"
+               OpName %func_F_i "func_F_i"
+               OpName %pointer "pointer"
+               OpName %main "main"
+               OpName %F "F"
+               OpMemberDecorate %str 0 Offset 0
+      %float = OpTypeFloat 32
+    %v4float = OpTypeVector %float 4
+        %str = OpTypeStruct %v4float
+%_ptr_Function_str = OpTypePointer Function %str
+          %1 = OpTypeFunction %v4float %_ptr_Function_str
+       %uint = OpTypeInt 32 0
+     %uint_0 = OpConstant %uint 0
+%_ptr_Function_v4float = OpTypePointer Function %v4float
+       %void = OpTypeVoid
+         %15 = OpTypeFunction %void
+         %20 = OpConstantNull %str
+   %func_F_i = OpFunction %v4float None %1
+    %pointer = OpFunctionParameter %_ptr_Function_str
+          %8 = OpLabel
+         %13 = OpAccessChain %_ptr_Function_v4float %pointer %uint_0
+         %14 = OpLoad %v4float %13
+               OpReturnValue %14
+               OpFunctionEnd
+       %main = OpFunction %void None %15
+         %18 = OpLabel
+          %F = OpVariable %_ptr_Function_str Function %20
+         %21 = OpFunctionCall %v4float %func_F_i %F
+               OpReturn
+               OpFunctionEnd

test/tint/ptr_ref/load/param/function/vec4_f32_in_struct.wgsl.expected.wgsl

@@ -0,0 +1,15 @@
+enable chromium_experimental_full_ptr_parameters;
+
+struct str {
+  i : vec4<f32>,
+}
+
+fn func(pointer : ptr<function, vec4<f32>>) -> vec4<f32> {
+  return *(pointer);
+}
+
+@compute @workgroup_size(1)
+fn main() {
+  var F : str;
+  let r = func(&(F.i));
+}

test/tint/ptr_ref/load/param/private/array_in_struct.wgsl

@@ -0,0 +1,16 @@
+enable chromium_experimental_full_ptr_parameters;
+
+struct str {
+  arr : array<i32, 4>,
+};
+
+fn func(pointer : ptr<private, array<i32, 4>>) -> array<i32, 4> {
+  return *pointer;
+}
+
+var<private> P : str;
+
+@compute @workgroup_size(1)
+fn main() {
+  let r = func(&P.arr);
+}

test/tint/ptr_ref/load/param/private/array_in_struct.wgsl.expected.dxc.hlsl

@@ -0,0 +1,16 @@
+struct str {
+  int arr[4];
+};
+
+typedef int func_ret[4];
+func_ret func(inout int pointer[4]) {
+  return pointer;
+}
+
+static str P = (str)0;
+
+[numthreads(1, 1, 1)]
+void main() {
+  const int r[4] = func(P.arr);
+  return;
+}

test/tint/ptr_ref/load/param/private/array_in_struct.wgsl.expected.fxc.hlsl

@@ -0,0 +1,16 @@
+struct str {
+  int arr[4];
+};
+
+typedef int func_ret[4];
+func_ret func(inout int pointer[4]) {
+  return pointer;
+}
+
+static str P = (str)0;
+
+[numthreads(1, 1, 1)]
+void main() {
+  const int r[4] = func(P.arr);
+  return;
+}

test/tint/ptr_ref/load/param/private/array_in_struct.wgsl.expected.glsl

@@ -0,0 +1,20 @@
+#version 310 es
+
+struct str {
+  int arr[4];
+};
+
+int[4] func(inout int pointer[4]) {
+  return pointer;
+}
+
+str P = str(int[4](0, 0, 0, 0));
+void tint_symbol() {
+  int r[4] = func(P.arr);
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  tint_symbol();
+  return;
+}

test/tint/ptr_ref/load/param/private/array_in_struct.wgsl.expected.msl

@@ -0,0 +1,30 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+template<typename T, size_t N>
+struct tint_array {
+    const constant T& operator[](size_t i) const constant { return elements[i]; }
+    device T& operator[](size_t i) device { return elements[i]; }
+    const device T& operator[](size_t i) const device { return elements[i]; }
+    thread T& operator[](size_t i) thread { return elements[i]; }
+    const thread T& operator[](size_t i) const thread { return elements[i]; }
+    threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+    const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+    T elements[N];
+};
+
+struct str {
+  tint_array<int, 4> arr;
+};
+
+tint_array<int, 4> func(thread tint_array<int, 4>* const pointer) {
+  return *(pointer);
+}
+
+kernel void tint_symbol() {
+  thread str tint_symbol_1 = {};
+  tint_array<int, 4> const r = func(&(tint_symbol_1.arr));
+  return;
+}
+

test/tint/ptr_ref/load/param/private/array_in_struct.wgsl.expected.spvasm

@@ -0,0 +1,42 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 24
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %str "str"
+               OpMemberName %str 0 "arr"
+               OpName %P "P"
+               OpName %func_F_arr "func_F_arr"
+               OpName %pointer "pointer"
+               OpName %main "main"
+               OpMemberDecorate %str 0 Offset 0
+               OpDecorate %_arr_int_uint_4 ArrayStride 4
+        %int = OpTypeInt 32 1
+       %uint = OpTypeInt 32 0
+     %uint_4 = OpConstant %uint 4
+%_arr_int_uint_4 = OpTypeArray %int %uint_4
+        %str = OpTypeStruct %_arr_int_uint_4
+%_ptr_Private_str = OpTypePointer Private %str
+          %8 = OpConstantNull %str
+          %P = OpVariable %_ptr_Private_str Private %8
+          %9 = OpTypeFunction %_arr_int_uint_4 %_ptr_Private_str
+     %uint_0 = OpConstant %uint 0
+%_ptr_Private__arr_int_uint_4 = OpTypePointer Private %_arr_int_uint_4
+       %void = OpTypeVoid
+         %18 = OpTypeFunction %void
+ %func_F_arr = OpFunction %_arr_int_uint_4 None %9
+    %pointer = OpFunctionParameter %_ptr_Private_str
+         %12 = OpLabel
+         %16 = OpAccessChain %_ptr_Private__arr_int_uint_4 %pointer %uint_0
+         %17 = OpLoad %_arr_int_uint_4 %16
+               OpReturnValue %17
+               OpFunctionEnd
+       %main = OpFunction %void None %18
+         %21 = OpLabel
+         %22 = OpFunctionCall %_arr_int_uint_4 %func_F_arr %P
+               OpReturn
+               OpFunctionEnd

test/tint/ptr_ref/load/param/private/array_in_struct.wgsl.expected.wgsl

@@ -0,0 +1,16 @@
+enable chromium_experimental_full_ptr_parameters;
+
+struct str {
+  arr : array<i32, 4>,
+}
+
+fn func(pointer : ptr<private, array<i32, 4>>) -> array<i32, 4> {
+  return *(pointer);
+}
+
+var<private> P : str;
+
+@compute @workgroup_size(1)
+fn main() {
+  let r = func(&(P.arr));
+}

test/tint/ptr_ref/load/param/private/i32.wgsl

@@ -0,0 +1,10 @@
+fn func(pointer : ptr<private, i32>) -> i32 {
+  return *pointer;
+}
+
+var<private> P : i32;
+
+@compute @workgroup_size(1)
+fn main() {
+  let r = func(&P);
+}

test/tint/ptr_ref/load/param/private/i32.wgsl.expected.dxc.hlsl

@@ -0,0 +1,11 @@
+int func(inout int pointer) {
+  return pointer;
+}
+
+static int P = 0;
+
+[numthreads(1, 1, 1)]
+void main() {
+  const int r = func(P);
+  return;
+}

test/tint/ptr_ref/load/param/private/i32.wgsl.expected.fxc.hlsl

@@ -0,0 +1,11 @@
+int func(inout int pointer) {
+  return pointer;
+}
+
+static int P = 0;
+
+[numthreads(1, 1, 1)]
+void main() {
+  const int r = func(P);
+  return;
+}

test/tint/ptr_ref/load/param/private/i32.wgsl.expected.glsl

@@ -0,0 +1,16 @@
+#version 310 es
+
+int func(inout int pointer) {
+  return pointer;
+}
+
+int P = 0;
+void tint_symbol() {
+  int r = func(P);
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  tint_symbol();
+  return;
+}

test/tint/ptr_ref/load/param/private/i32.wgsl.expected.msl

@@ -0,0 +1,13 @@
+#include <metal_stdlib>
+
+using namespace metal;
+int func(thread int* const pointer) {
+  return *(pointer);
+}
+
+kernel void tint_symbol() {
+  thread int tint_symbol_1 = 0;
+  int const r = func(&(tint_symbol_1));
+  return;
+}
+

test/tint/ptr_ref/load/param/private/i32.wgsl.expected.spvasm

@@ -0,0 +1,31 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 17
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %P "P"
+               OpName %func "func"
+               OpName %pointer "pointer"
+               OpName %main "main"
+        %int = OpTypeInt 32 1
+%_ptr_Private_int = OpTypePointer Private %int
+          %4 = OpConstantNull %int
+          %P = OpVariable %_ptr_Private_int Private %4
+          %5 = OpTypeFunction %int %_ptr_Private_int
+       %void = OpTypeVoid
+         %11 = OpTypeFunction %void
+       %func = OpFunction %int None %5
+    %pointer = OpFunctionParameter %_ptr_Private_int
+          %8 = OpLabel
+         %10 = OpLoad %int %pointer
+               OpReturnValue %10
+               OpFunctionEnd
+       %main = OpFunction %void None %11
+         %14 = OpLabel
+         %15 = OpFunctionCall %int %func %P
+               OpReturn
+               OpFunctionEnd

test/tint/ptr_ref/load/param/private/i32.wgsl.expected.wgsl

@@ -0,0 +1,10 @@
+fn func(pointer : ptr<private, i32>) -> i32 {
+  return *(pointer);
+}
+
+var<private> P : i32;
+
+@compute @workgroup_size(1)
+fn main() {
+  let r = func(&(P));
+}

test/tint/ptr_ref/load/param/private/i32_in_struct.wgsl

@@ -0,0 +1,16 @@
+enable chromium_experimental_full_ptr_parameters;
+
+struct str {
+  i : i32,
+};
+
+fn func(pointer : ptr<private, i32>) -> i32 {
+  return *pointer;
+}
+
+var<private> P : str;
+
+@compute @workgroup_size(1)
+fn main() {
+  let r = func(&P.i);
+}

test/tint/ptr_ref/load/param/private/i32_in_struct.wgsl.expected.dxc.hlsl

@@ -0,0 +1,15 @@
+struct str {
+  int i;
+};
+
+int func(inout int pointer) {
+  return pointer;
+}
+
+static str P = (str)0;
+
+[numthreads(1, 1, 1)]
+void main() {
+  const int r = func(P.i);
+  return;
+}

test/tint/ptr_ref/load/param/private/i32_in_struct.wgsl.expected.fxc.hlsl

@@ -0,0 +1,15 @@
+struct str {
+  int i;
+};
+
+int func(inout int pointer) {
+  return pointer;
+}
+
+static str P = (str)0;
+
+[numthreads(1, 1, 1)]
+void main() {
+  const int r = func(P.i);
+  return;
+}

test/tint/ptr_ref/load/param/private/i32_in_struct.wgsl.expected.glsl

@@ -0,0 +1,20 @@
+#version 310 es
+
+struct str {
+  int i;
+};
+
+int func(inout int pointer) {
+  return pointer;
+}
+
+str P = str(0);
+void tint_symbol() {
+  int r = func(P.i);
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  tint_symbol();
+  return;
+}

test/tint/ptr_ref/load/param/private/i32_in_struct.wgsl.expected.msl

@@ -0,0 +1,17 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct str {
+  int i;
+};
+
+int func(thread int* const pointer) {
+  return *(pointer);
+}
+
+kernel void tint_symbol() {
+  thread str tint_symbol_1 = {};
+  int const r = func(&(tint_symbol_1.i));
+  return;
+}
+

test/tint/ptr_ref/load/param/private/i32_in_struct.wgsl.expected.spvasm

@@ -0,0 +1,39 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 22
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %str "str"
+               OpMemberName %str 0 "i"
+               OpName %P "P"
+               OpName %func_F_i "func_F_i"
+               OpName %pointer "pointer"
+               OpName %main "main"
+               OpMemberDecorate %str 0 Offset 0
+        %int = OpTypeInt 32 1
+        %str = OpTypeStruct %int
+%_ptr_Private_str = OpTypePointer Private %str
+          %5 = OpConstantNull %str
+          %P = OpVariable %_ptr_Private_str Private %5
+          %6 = OpTypeFunction %int %_ptr_Private_str
+       %uint = OpTypeInt 32 0
+     %uint_0 = OpConstant %uint 0
+%_ptr_Private_int = OpTypePointer Private %int
+       %void = OpTypeVoid
+         %16 = OpTypeFunction %void
+   %func_F_i = OpFunction %int None %6
+    %pointer = OpFunctionParameter %_ptr_Private_str
+          %9 = OpLabel
+         %14 = OpAccessChain %_ptr_Private_int %pointer %uint_0
+         %15 = OpLoad %int %14
+               OpReturnValue %15
+               OpFunctionEnd
+       %main = OpFunction %void None %16
+         %19 = OpLabel
+         %20 = OpFunctionCall %int %func_F_i %P
+               OpReturn
+               OpFunctionEnd

test/tint/ptr_ref/load/param/private/i32_in_struct.wgsl.expected.wgsl

@@ -0,0 +1,16 @@
+enable chromium_experimental_full_ptr_parameters;
+
+struct str {
+  i : i32,
+}
+
+fn func(pointer : ptr<private, i32>) -> i32 {
+  return *(pointer);
+}
+
+var<private> P : str;
+
+@compute @workgroup_size(1)
+fn main() {
+  let r = func(&(P.i));
+}

test/tint/ptr_ref/load/param/private/struct_in_array.wgsl

@@ -0,0 +1,16 @@
+enable chromium_experimental_full_ptr_parameters;
+
+struct str {
+  i : i32,
+};
+
+fn func(pointer : ptr<private, str>) -> str {
+  return *pointer;
+}
+
+var<private> P : array<str, 4>;
+
+@compute @workgroup_size(1)
+fn main() {
+  let r = func(&P[2]);
+}

test/tint/ptr_ref/load/param/private/struct_in_array.wgsl.expected.dxc.hlsl

@@ -0,0 +1,15 @@
+struct str {
+  int i;
+};
+
+str func(inout str pointer) {
+  return pointer;
+}
+
+static str P[4] = (str[4])0;
+
+[numthreads(1, 1, 1)]
+void main() {
+  const str r = func(P[2]);
+  return;
+}

test/tint/ptr_ref/load/param/private/struct_in_array.wgsl.expected.fxc.hlsl

@@ -0,0 +1,15 @@
+struct str {
+  int i;
+};
+
+str func(inout str pointer) {
+  return pointer;
+}
+
+static str P[4] = (str[4])0;
+
+[numthreads(1, 1, 1)]
+void main() {
+  const str r = func(P[2]);
+  return;
+}

test/tint/ptr_ref/load/param/private/struct_in_array.wgsl.expected.glsl

@@ -0,0 +1,20 @@
+#version 310 es
+
+struct str {
+  int i;
+};
+
+str func(inout str pointer) {
+  return pointer;
+}
+
+str P[4] = str[4](str(0), str(0), str(0), str(0));
+void tint_symbol() {
+  str r = func(P[2]);
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  tint_symbol();
+  return;
+}

test/tint/ptr_ref/load/param/private/struct_in_array.wgsl.expected.msl

@@ -0,0 +1,30 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+template<typename T, size_t N>
+struct tint_array {
+    const constant T& operator[](size_t i) const constant { return elements[i]; }
+    device T& operator[](size_t i) device { return elements[i]; }
+    const device T& operator[](size_t i) const device { return elements[i]; }
+    thread T& operator[](size_t i) thread { return elements[i]; }
+    const thread T& operator[](size_t i) const thread { return elements[i]; }
+    threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+    const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+    T elements[N];
+};
+
+struct str {
+  int i;
+};
+
+str func(thread str* const pointer) {
+  return *(pointer);
+}
+
+kernel void tint_symbol() {
+  thread tint_array<str, 4> tint_symbol_1 = {};
+  str const r = func(&(tint_symbol_1[2]));
+  return;
+}
+

test/tint/ptr_ref/load/param/private/struct_in_array.wgsl.expected.spvasm

@@ -0,0 +1,50 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 30
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %str "str"
+               OpMemberName %str 0 "i"
+               OpName %P "P"
+               OpName %func_F_X "func_F_X"
+               OpName %pointer_base "pointer_base"
+               OpName %pointer_indices "pointer_indices"
+               OpName %main "main"
+               OpMemberDecorate %str 0 Offset 0
+               OpDecorate %_arr_str_uint_4 ArrayStride 4
+               OpDecorate %_arr_uint_uint_1 ArrayStride 4
+        %int = OpTypeInt 32 1
+        %str = OpTypeStruct %int
+       %uint = OpTypeInt 32 0
+     %uint_4 = OpConstant %uint 4
+%_arr_str_uint_4 = OpTypeArray %str %uint_4
+%_ptr_Private__arr_str_uint_4 = OpTypePointer Private %_arr_str_uint_4
+          %8 = OpConstantNull %_arr_str_uint_4
+          %P = OpVariable %_ptr_Private__arr_str_uint_4 Private %8
+     %uint_1 = OpConstant %uint 1
+%_arr_uint_uint_1 = OpTypeArray %uint %uint_1
+          %9 = OpTypeFunction %str %_ptr_Private__arr_str_uint_4 %_arr_uint_uint_1
+         %17 = OpConstantNull %int
+%_ptr_Private_str = OpTypePointer Private %str
+       %void = OpTypeVoid
+         %22 = OpTypeFunction %void
+     %uint_2 = OpConstant %uint 2
+         %29 = OpConstantComposite %_arr_uint_uint_1 %uint_2
+   %func_F_X = OpFunction %str None %9
+%pointer_base = OpFunctionParameter %_ptr_Private__arr_str_uint_4
+%pointer_indices = OpFunctionParameter %_arr_uint_uint_1
+         %15 = OpLabel
+         %18 = OpCompositeExtract %uint %pointer_indices 0
+         %20 = OpAccessChain %_ptr_Private_str %pointer_base %18
+         %21 = OpLoad %str %20
+               OpReturnValue %21
+               OpFunctionEnd
+       %main = OpFunction %void None %22
+         %25 = OpLabel
+         %26 = OpFunctionCall %str %func_F_X %P %29
+               OpReturn
+               OpFunctionEnd

test/tint/ptr_ref/load/param/private/struct_in_array.wgsl.expected.wgsl

@@ -0,0 +1,16 @@
+enable chromium_experimental_full_ptr_parameters;
+
+struct str {
+  i : i32,
+}
+
+fn func(pointer : ptr<private, str>) -> str {
+  return *(pointer);
+}
+
+var<private> P : array<str, 4>;
+
+@compute @workgroup_size(1)
+fn main() {
+  let r = func(&(P[2]));
+}

test/tint/ptr_ref/load/param/private/vec2_f32_in_mat2x2.wgsl

@@ -0,0 +1,12 @@
+enable chromium_experimental_full_ptr_parameters;
+
+fn func(pointer : ptr<private, vec2<f32>>) -> vec2<f32> {
+  return *pointer;
+}
+
+var<private> P : mat2x2<f32>;
+
+@compute @workgroup_size(1)
+fn main() {
+  let r = func(&P[1]);
+}

test/tint/ptr_ref/load/param/private/vec2_f32_in_mat2x2.wgsl.expected.dxc.hlsl

@@ -0,0 +1,11 @@
+float2 func(inout float2 pointer) {
+  return pointer;
+}
+
+static float2x2 P = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+
+[numthreads(1, 1, 1)]
+void main() {
+  const float2 r = func(P[1]);
+  return;
+}

test/tint/ptr_ref/load/param/private/vec2_f32_in_mat2x2.wgsl.expected.fxc.hlsl

@@ -0,0 +1,11 @@
+float2 func(inout float2 pointer) {
+  return pointer;
+}
+
+static float2x2 P = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+
+[numthreads(1, 1, 1)]
+void main() {
+  const float2 r = func(P[1]);
+  return;
+}

test/tint/ptr_ref/load/param/private/vec2_f32_in_mat2x2.wgsl.expected.glsl

@@ -0,0 +1,16 @@
+#version 310 es
+
+vec2 func(inout vec2 pointer) {
+  return pointer;
+}
+
+mat2 P = mat2(0.0f, 0.0f, 0.0f, 0.0f);
+void tint_symbol() {
+  vec2 r = func(P[1]);
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  tint_symbol();
+  return;
+}