spirv-reader: switch to HLSL-style pipeline IO

- When storing to sample_mask output, write to the 0th element
- Only make a return struct if it has members
- Adjust type signedness coercion when loading special builtins.
- Adapt tests

- Update expectations for end-to-end tests

- Handle sample_mask with stride
  Input variables normally don't have layout. But they can have it
  up through SPIR-V 1.4.
  Handle this case in the SPIR-V reader, by seeing through the
  intermediate alias type created for the strided array type.

Bug: tint:508
Change-Id: I0f19dc1305d3f250dbbc0698a602288c34245274
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/54743
Auto-Submit: David Neto <dneto@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Reviewed-by: James Price <jrprice@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>

test/access/let/matrix.spvasm.expected.hlsl

@@ -1,5 +1,10 @@
-[numthreads(1, 1, 1)]
-void main() {
+void main_1() {
   const float x_24 = float3x3(float3(1.0f, 2.0f, 3.0f), float3(4.0f, 5.0f, 6.0f), float3(7.0f, 8.0f, 9.0f))[1u].y;
   return;
 }
+
+[numthreads(1, 1, 1)]
+void main() {
+  main_1();
+  return;
+}

test/access/let/matrix.spvasm.expected.msl

@@ -1,8 +1,13 @@
 #include <metal_stdlib>
 
 using namespace metal;
-kernel void tint_symbol() {
+void main_1() {
   float const x_24 = float3x3(float3(1.0f, 2.0f, 3.0f), float3(4.0f, 5.0f, 6.0f), float3(7.0f, 8.0f, 9.0f))[1u].y;
   return;
 }
 
+kernel void tint_symbol() {
+  main_1();
+  return;
+}
+

test/access/let/matrix.spvasm.expected.spvasm

@@ -1,12 +1,13 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 25
+; Bound: 28
 ; Schema: 0
                OpCapability Shader
                OpMemoryModel Logical GLSL450
                OpEntryPoint GLCompute %main "main"
                OpExecutionMode %main LocalSize 1 1 1
+               OpName %main_1 "main_1"
                OpName %main "main"
        %void = OpTypeVoid
           %1 = OpTypeFunction %void
@@ -28,9 +29,14 @@
          %20 = OpConstantComposite %mat3v3float %11 %15 %19
        %uint = OpTypeInt 32 0
      %uint_1 = OpConstant %uint 1
-       %main = OpFunction %void None %1
+     %main_1 = OpFunction %void None %1
           %4 = OpLabel
          %23 = OpCompositeExtract %v3float %20 1
          %24 = OpCompositeExtract %float %23 1
                OpReturn
                OpFunctionEnd
+       %main = OpFunction %void None %1
+         %26 = OpLabel
+         %27 = OpFunctionCall %void %main_1
+               OpReturn
+               OpFunctionEnd

test/access/let/matrix.spvasm.expected.wgsl

@@ -1,5 +1,9 @@
-[[stage(compute)]]
-fn main() {
+fn main_1() {
   let x_24 : f32 = mat3x3<f32>(vec3<f32>(1.0, 2.0, 3.0), vec3<f32>(4.0, 5.0, 6.0), vec3<f32>(7.0, 8.0, 9.0))[1u].y;
   return;
 }
+
+[[stage(compute)]]
+fn main() {
+  main_1();
+}

test/access/let/vector.spvasm.expected.hlsl

@@ -1,7 +1,12 @@
-[numthreads(1, 1, 1)]
-void main() {
+void main_1() {
   const float x_11 = float3(1.0f, 2.0f, 3.0f).y;
   const float2 x_13 = float2(float3(1.0f, 2.0f, 3.0f).x, float3(1.0f, 2.0f, 3.0f).z);
   const float3 x_14 = float3(float3(1.0f, 2.0f, 3.0f).x, float3(1.0f, 2.0f, 3.0f).z, float3(1.0f, 2.0f, 3.0f).y);
   return;
 }
+
+[numthreads(1, 1, 1)]
+void main() {
+  main_1();
+  return;
+}

test/access/let/vector.spvasm.expected.msl

@@ -1,10 +1,15 @@
 #include <metal_stdlib>
 
 using namespace metal;
-kernel void tint_symbol() {
+void main_1() {
   float const x_11 = float3(1.0f, 2.0f, 3.0f).y;
   float2 const x_13 = float2(float3(1.0f, 2.0f, 3.0f).x, float3(1.0f, 2.0f, 3.0f).z);
   float3 const x_14 = float3(float3(1.0f, 2.0f, 3.0f).x, float3(1.0f, 2.0f, 3.0f).z, float3(1.0f, 2.0f, 3.0f).y);
   return;
 }
 
+kernel void tint_symbol() {
+  main_1();
+  return;
+}
+

test/access/let/vector.spvasm.expected.spvasm

@@ -1,12 +1,13 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 20
+; Bound: 23
 ; Schema: 0
                OpCapability Shader
                OpMemoryModel Logical GLSL450
                OpEntryPoint GLCompute %main "main"
                OpExecutionMode %main LocalSize 1 1 1
+               OpName %main_1 "main_1"
                OpName %main "main"
        %void = OpTypeVoid
           %1 = OpTypeFunction %void
@@ -17,7 +18,7 @@
     %float_3 = OpConstant %float 3
          %10 = OpConstantComposite %v3float %float_1 %float_2 %float_3
     %v2float = OpTypeVector %float 2
-       %main = OpFunction %void None %1
+     %main_1 = OpFunction %void None %1
           %4 = OpLabel
          %11 = OpCompositeExtract %float %10 1
          %13 = OpCompositeExtract %float %10 0
@@ -29,3 +30,8 @@
          %19 = OpCompositeConstruct %v3float %16 %17 %18
                OpReturn
                OpFunctionEnd
+       %main = OpFunction %void None %1
+         %21 = OpLabel
+         %22 = OpFunctionCall %void %main_1
+               OpReturn
+               OpFunctionEnd

test/access/let/vector.spvasm.expected.wgsl

@@ -1,7 +1,11 @@
-[[stage(compute)]]
-fn main() {
+fn main_1() {
   let x_11 : f32 = vec3<f32>(1.0, 2.0, 3.0).y;
   let x_13 : vec2<f32> = vec2<f32>(vec3<f32>(1.0, 2.0, 3.0).x, vec3<f32>(1.0, 2.0, 3.0).z);
   let x_14 : vec3<f32> = vec3<f32>(vec3<f32>(1.0, 2.0, 3.0).x, vec3<f32>(1.0, 2.0, 3.0).z, vec3<f32>(1.0, 2.0, 3.0).y);
   return;
 }
+
+[[stage(compute)]]
+fn main() {
+  main_1();
+}

test/access/var/matrix.spvasm.expected.hlsl

@@ -1,7 +1,12 @@
-[numthreads(1, 1, 1)]
-void main() {
+void main_1() {
   float3x3 m = float3x3(float3(0.0f, 0.0f, 0.0f), float3(0.0f, 0.0f, 0.0f), float3(0.0f, 0.0f, 0.0f));
   const float3 x_15 = m[1];
   const float x_16 = x_15.y;
   return;
 }
+
+[numthreads(1, 1, 1)]
+void main() {
+  main_1();
+  return;
+}

test/access/var/matrix.spvasm.expected.msl

@@ -1,10 +1,15 @@
 #include <metal_stdlib>
 
 using namespace metal;
-kernel void tint_symbol() {
+void main_1() {
   float3x3 m = float3x3(float3(0.0f, 0.0f, 0.0f), float3(0.0f, 0.0f, 0.0f), float3(0.0f, 0.0f, 0.0f));
   float3 const x_15 = m[1];
   float const x_16 = x_15.y;
   return;
 }
 
+kernel void tint_symbol() {
+  main_1();
+  return;
+}
+

test/access/var/matrix.spvasm.expected.spvasm

@@ -1,14 +1,15 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 20
+; Bound: 23
 ; Schema: 0
                OpCapability Shader
                OpMemoryModel Logical GLSL450
                OpEntryPoint GLCompute %main "main"
                OpExecutionMode %main LocalSize 1 1 1
-               OpName %main "main"
+               OpName %main_1 "main_1"
                OpName %m "m"
+               OpName %main "main"
        %void = OpTypeVoid
           %1 = OpTypeFunction %void
       %float = OpTypeFloat 32
@@ -22,7 +23,7 @@
         %int = OpTypeInt 32 1
       %int_1 = OpConstant %int 1
 %_ptr_Function_v3float = OpTypePointer Function %v3float
-       %main = OpFunction %void None %1
+     %main_1 = OpFunction %void None %1
           %4 = OpLabel
           %m = OpVariable %_ptr_Function_mat3v3float Function %13
                OpStore %m %10
@@ -31,3 +32,8 @@
          %19 = OpCompositeExtract %float %18 1
                OpReturn
                OpFunctionEnd
+       %main = OpFunction %void None %1
+         %21 = OpLabel
+         %22 = OpFunctionCall %void %main_1
+               OpReturn
+               OpFunctionEnd

test/access/var/matrix.spvasm.expected.wgsl

@@ -1,7 +1,11 @@
-[[stage(compute)]]
-fn main() {
+fn main_1() {
   var m : mat3x3<f32> = mat3x3<f32>(vec3<f32>(0.0, 0.0, 0.0), vec3<f32>(0.0, 0.0, 0.0), vec3<f32>(0.0, 0.0, 0.0));
   let x_15 : vec3<f32> = m[1];
   let x_16 : f32 = x_15.y;
   return;
 }
+
+[[stage(compute)]]
+fn main() {
+  main_1();
+}

test/access/var/vector.spvasm.expected.hlsl

@@ -1,5 +1,4 @@
-[numthreads(1, 1, 1)]
-void main() {
+void main_1() {
   float3 v = float3(0.0f, 0.0f, 0.0f);
   const float x_14 = v.y;
   const float3 x_16 = v;
@@ -8,3 +7,9 @@ void main() {
   const float3 x_19 = float3(x_18.x, x_18.z, x_18.y);
   return;
 }
+
+[numthreads(1, 1, 1)]
+void main() {
+  main_1();
+  return;
+}

test/access/var/vector.spvasm.expected.msl

@@ -1,7 +1,7 @@
 #include <metal_stdlib>
 
 using namespace metal;
-kernel void tint_symbol() {
+void main_1() {
   float3 v = float3(0.0f, 0.0f, 0.0f);
   float const x_14 = v.y;
   float3 const x_16 = v;
@@ -11,3 +11,8 @@ kernel void tint_symbol() {
   return;
 }
 
+kernel void tint_symbol() {
+  main_1();
+  return;
+}
+

test/access/var/vector.spvasm.expected.spvasm

@@ -1,14 +1,15 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 27
+; Bound: 30
 ; Schema: 0
                OpCapability Shader
                OpMemoryModel Logical GLSL450
                OpEntryPoint GLCompute %main "main"
                OpExecutionMode %main LocalSize 1 1 1
-               OpName %main "main"
+               OpName %main_1 "main_1"
                OpName %v "v"
+               OpName %main "main"
        %void = OpTypeVoid
           %1 = OpTypeFunction %void
       %float = OpTypeFloat 32
@@ -21,7 +22,7 @@
      %uint_1 = OpConstant %uint 1
 %_ptr_Function_float = OpTypePointer Function %float
     %v2float = OpTypeVector %float 2
-       %main = OpFunction %void None %1
+     %main_1 = OpFunction %void None %1
           %4 = OpLabel
           %v = OpVariable %_ptr_Function_v3float Function %11
                OpStore %v %8
@@ -38,3 +39,8 @@
          %26 = OpCompositeConstruct %v3float %23 %24 %25
                OpReturn
                OpFunctionEnd
+       %main = OpFunction %void None %1
+         %28 = OpLabel
+         %29 = OpFunctionCall %void %main_1
+               OpReturn
+               OpFunctionEnd

test/access/var/vector.spvasm.expected.wgsl

@@ -1,5 +1,4 @@
-[[stage(compute)]]
-fn main() {
+fn main_1() {
   var v : vec3<f32> = vec3<f32>(0.0, 0.0, 0.0);
   let x_14 : f32 = v.y;
   let x_16 : vec3<f32> = v;
@@ -8,3 +7,8 @@ fn main() {
   let x_19 : vec3<f32> = vec3<f32>(x_18.x, x_18.z, x_18.y);
   return;
 }
+
+[[stage(compute)]]
+fn main() {
+  main_1();
+}

test/array/assign_to_function_var.wgsl.expected.msl

@@ -1 +1,53 @@
-SKIP: crbug.com/tint/649
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_padded_array_element {
+  /* 0x0000 */ int el;
+  /* 0x0004 */ int8_t tint_pad_0[12];
+};
+struct tint_array_wrapper {
+  /* 0x0000 */ tint_padded_array_element arr[4];
+};
+struct S {
+  /* 0x0000 */ tint_array_wrapper arr;
+};
+struct tint_array_wrapper_3 {
+  int arr[2];
+};
+struct tint_array_wrapper_2 {
+  tint_array_wrapper_3 arr[3];
+};
+struct tint_array_wrapper_1 {
+  tint_array_wrapper_2 arr[4];
+};
+
+tint_array_wrapper ret_arr() {
+  tint_array_wrapper const tint_symbol = {.arr={}};
+  return tint_symbol;
+}
+
+S ret_struct_arr() {
+  S const tint_symbol_1 = {};
+  return tint_symbol_1;
+}
+
+void foo(constant S& src_uniform, device S& src_storage, tint_array_wrapper src_param, thread tint_array_wrapper* const tint_symbol_3, threadgroup tint_array_wrapper* const tint_symbol_4) {
+  tint_array_wrapper src_function = {};
+  tint_array_wrapper dst = {};
+  tint_array_wrapper const tint_symbol_2 = {.arr={{.el=1}, {.el=2}, {.el=3}, {.el=3}}};
+  dst = tint_symbol_2;
+  dst = src_param;
+  dst = ret_arr();
+  tint_array_wrapper const src_let = {.arr={}};
+  dst = src_let;
+  dst = src_function;
+  dst = *(tint_symbol_3);
+  dst = *(tint_symbol_4);
+  dst = ret_struct_arr().arr;
+  dst = src_uniform.arr;
+  dst = src_storage.arr;
+  tint_array_wrapper_1 dst_nested = {};
+  tint_array_wrapper_1 src_nested = {};
+  dst_nested = src_nested;
+}
+

test/array/assign_to_private_var.wgsl.expected.msl

@@ -1 +1,51 @@
-SKIP: crbug.com/tint/649
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_padded_array_element {
+  /* 0x0000 */ int el;
+  /* 0x0004 */ int8_t tint_pad_0[12];
+};
+struct tint_array_wrapper {
+  /* 0x0000 */ tint_padded_array_element arr[4];
+};
+struct S {
+  /* 0x0000 */ tint_array_wrapper arr;
+};
+struct tint_array_wrapper_3 {
+  int arr[2];
+};
+struct tint_array_wrapper_2 {
+  tint_array_wrapper_3 arr[3];
+};
+struct tint_array_wrapper_1 {
+  tint_array_wrapper_2 arr[4];
+};
+
+tint_array_wrapper ret_arr() {
+  tint_array_wrapper const tint_symbol = {.arr={}};
+  return tint_symbol;
+}
+
+S ret_struct_arr() {
+  S const tint_symbol_1 = {};
+  return tint_symbol_1;
+}
+
+void foo(constant S& src_uniform, device S& src_storage, tint_array_wrapper src_param, thread tint_array_wrapper* const tint_symbol_3, thread tint_array_wrapper* const tint_symbol_4, threadgroup tint_array_wrapper* const tint_symbol_5, thread tint_array_wrapper_1* const tint_symbol_6) {
+  tint_array_wrapper src_function = {};
+  tint_array_wrapper const tint_symbol_2 = {.arr={{.el=1}, {.el=2}, {.el=3}, {.el=3}}};
+  *(tint_symbol_3) = tint_symbol_2;
+  *(tint_symbol_3) = src_param;
+  *(tint_symbol_3) = ret_arr();
+  tint_array_wrapper const src_let = {.arr={}};
+  *(tint_symbol_3) = src_let;
+  *(tint_symbol_3) = src_function;
+  *(tint_symbol_3) = *(tint_symbol_4);
+  *(tint_symbol_3) = *(tint_symbol_5);
+  *(tint_symbol_3) = ret_struct_arr().arr;
+  *(tint_symbol_3) = src_uniform.arr;
+  *(tint_symbol_3) = src_storage.arr;
+  tint_array_wrapper_1 src_nested = {};
+  *(tint_symbol_6) = src_nested;
+}
+

test/array/assign_to_storage_var.wgsl.expected.msl

@@ -1 +1,54 @@
-SKIP: crbug.com/tint/649
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_padded_array_element {
+  /* 0x0000 */ int el;
+  /* 0x0004 */ int8_t tint_pad_0[12];
+};
+struct tint_array_wrapper {
+  /* 0x0000 */ tint_padded_array_element arr[4];
+};
+struct S {
+  /* 0x0000 */ tint_array_wrapper arr;
+};
+struct tint_array_wrapper_3 {
+  /* 0x0000 */ int arr[2];
+};
+struct tint_array_wrapper_2 {
+  /* 0x0000 */ tint_array_wrapper_3 arr[3];
+};
+struct tint_array_wrapper_1 {
+  /* 0x0000 */ tint_array_wrapper_2 arr[4];
+};
+struct S_nested {
+  /* 0x0000 */ tint_array_wrapper_1 arr;
+};
+
+tint_array_wrapper ret_arr() {
+  tint_array_wrapper const tint_symbol = {.arr={}};
+  return tint_symbol;
+}
+
+S ret_struct_arr() {
+  S const tint_symbol_1 = {};
+  return tint_symbol_1;
+}
+
+void foo(constant S& src_uniform, device S& dst, device S& src_storage, device S_nested& dst_nested, tint_array_wrapper src_param, thread tint_array_wrapper* const tint_symbol_3, threadgroup tint_array_wrapper* const tint_symbol_4) {
+  tint_array_wrapper src_function = {};
+  tint_array_wrapper const tint_symbol_2 = {.arr={{.el=1}, {.el=2}, {.el=3}, {.el=3}}};
+  dst.arr = tint_symbol_2;
+  dst.arr = src_param;
+  dst.arr = ret_arr();
+  tint_array_wrapper const src_let = {.arr={}};
+  dst.arr = src_let;
+  dst.arr = src_function;
+  dst.arr = *(tint_symbol_3);
+  dst.arr = *(tint_symbol_4);
+  dst.arr = ret_struct_arr().arr;
+  dst.arr = src_uniform.arr;
+  dst.arr = src_storage.arr;
+  tint_array_wrapper_1 src_nested = {};
+  dst_nested.arr = src_nested;
+}
+

test/array/assign_to_workgroup_var.wgsl.expected.msl

@@ -1 +1,51 @@
-SKIP: crbug.com/tint/649
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_padded_array_element {
+  /* 0x0000 */ int el;
+  /* 0x0004 */ int8_t tint_pad_0[12];
+};
+struct tint_array_wrapper {
+  /* 0x0000 */ tint_padded_array_element arr[4];
+};
+struct S {
+  /* 0x0000 */ tint_array_wrapper arr;
+};
+struct tint_array_wrapper_3 {
+  int arr[2];
+};
+struct tint_array_wrapper_2 {
+  tint_array_wrapper_3 arr[3];
+};
+struct tint_array_wrapper_1 {
+  tint_array_wrapper_2 arr[4];
+};
+
+tint_array_wrapper ret_arr() {
+  tint_array_wrapper const tint_symbol = {.arr={}};
+  return tint_symbol;
+}
+
+S ret_struct_arr() {
+  S const tint_symbol_1 = {};
+  return tint_symbol_1;
+}
+
+void foo(constant S& src_uniform, device S& src_storage, tint_array_wrapper src_param, threadgroup tint_array_wrapper* const tint_symbol_3, thread tint_array_wrapper* const tint_symbol_4, threadgroup tint_array_wrapper* const tint_symbol_5, threadgroup tint_array_wrapper_1* const tint_symbol_6) {
+  tint_array_wrapper src_function = {};
+  tint_array_wrapper const tint_symbol_2 = {.arr={{.el=1}, {.el=2}, {.el=3}, {.el=3}}};
+  *(tint_symbol_3) = tint_symbol_2;
+  *(tint_symbol_3) = src_param;
+  *(tint_symbol_3) = ret_arr();
+  tint_array_wrapper const src_let = {.arr={}};
+  *(tint_symbol_3) = src_let;
+  *(tint_symbol_3) = src_function;
+  *(tint_symbol_3) = *(tint_symbol_4);
+  *(tint_symbol_3) = *(tint_symbol_5);
+  *(tint_symbol_3) = ret_struct_arr().arr;
+  *(tint_symbol_3) = src_uniform.arr;
+  *(tint_symbol_3) = src_storage.arr;
+  tint_array_wrapper_1 src_nested = {};
+  *(tint_symbol_6) = src_nested;
+}
+

test/bug/tint/413.spvasm.expected.hlsl

@@ -1,8 +1,7 @@
 Texture2D<uint4> Src : register(t0, space0);
 RWTexture2D<uint4> Dst : register(u1, space0);
 
-[numthreads(1, 1, 1)]
-void main() {
+void main_1() {
   uint4 srcValue = uint4(0u, 0u, 0u, 0u);
   const uint4 x_18 = Src.Load(int3(0, 0, 0));
   srcValue = x_18;
@@ -12,3 +11,9 @@ void main() {
   Dst[int2(0, 0)] = x_27;
   return;
 }
+
+[numthreads(1, 1, 1)]
+void main() {
+  main_1();
+  return;
+}

test/bug/tint/413.spvasm.expected.spvasm

@@ -1,7 +1,7 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 32
+; Bound: 35
 ; Schema: 0
                OpCapability Shader
                OpMemoryModel Logical GLSL450
@@ -9,8 +9,9 @@
                OpExecutionMode %main LocalSize 1 1 1
                OpName %Src "Src"
                OpName %Dst "Dst"
-               OpName %main "main"
+               OpName %main_1 "main_1"
                OpName %srcValue "srcValue"
+               OpName %main "main"
                OpDecorate %Src NonWritable
                OpDecorate %Src DescriptorSet 0
                OpDecorate %Src Binding 0
@@ -35,7 +36,7 @@
      %uint_0 = OpConstant %uint 0
 %_ptr_Function_uint = OpTypePointer Function %uint
       %int_1 = OpConstant %int 1
-       %main = OpFunction %void None %7
+     %main_1 = OpFunction %void None %7
          %10 = OpLabel
    %srcValue = OpVariable %_ptr_Function_v4uint Function %14
          %16 = OpLoad %3 %Src
@@ -52,3 +53,8 @@
                OpImageWrite %31 %20 %29
                OpReturn
                OpFunctionEnd
+       %main = OpFunction %void None %7
+         %33 = OpLabel
+         %34 = OpFunctionCall %void %main_1
+               OpReturn
+               OpFunctionEnd

test/bug/tint/413.spvasm.expected.wgsl

@@ -2,8 +2,7 @@
 
 [[group(0), binding(1)]] var Dst : texture_storage_2d<r32uint, write>;
 
-[[stage(compute)]]
-fn main() {
+fn main_1() {
   var srcValue : vec4<u32>;
   let x_18 : vec4<u32> = textureLoad(Src, vec2<i32>(0, 0));
   srcValue = x_18;
@@ -14,3 +13,8 @@ fn main() {
   textureStore(Dst, vec2<i32>(0, 0), x_27);
   return;
 }
+
+[[stage(compute)]]
+fn main() {
+  main_1();
+}

test/bug/tint/749.spvasm.expected.wgsl

@@ -9,11 +9,11 @@ struct buf0 {
 
 var<private> obj : QuicksortObject;
 
-[[builtin(position)]] var<in> gl_FragCoord : vec4<f32>;
+var<private> gl_FragCoord : vec4<f32>;
 
 [[group(0), binding(0)]] var<uniform> x_188 : buf0;
 
-[[location(0)]] var<out> x_GLF_color : vec4<f32>;
+var<private> x_GLF_color : vec4<f32>;
 
 fn swap_i1_i1_(i : ptr<function, i32>, j : ptr<function, i32>) {
   var temp : i32;
@@ -770,8 +770,7 @@ fn quicksort_() {
   return;
 }
 
-[[stage(fragment)]]
-fn main() {
+fn main_1() {
   var color : vec3<f32>;
   var i_2 : i32;
   var uv : vec2<f32>;
@@ -1492,3 +1491,15 @@ fn main() {
   *(x_208) = x_931;
   return;
 }
+
+struct main_out {
+  [[location(0)]]
+  x_GLF_color : vec4<f32>;
+};
+
+[[stage(fragment)]]
+fn main([[builtin(position)]] gl_FragCoord_param : vec4<f32>) -> main_out {
+  gl_FragCoord = gl_FragCoord_param;
+  main_1();
+  return main_out(x_GLF_color);
+}

test/bug/tint/757.wgsl.expected.msl

@@ -1,19 +1,34 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-bug/tint/757.wgsl:3:5 warning: use of deprecated language feature: [[offset]] has been replaced with [[size]] and [[align]]
-  [[offset(0)]] level : i32;
-    ^^^^^^
+using namespace metal;
+struct Constants {
+  /* 0x0000 */ int level;
+};
+struct Result {
+  /* 0x0000 */ float values[1];
+};
 
-bug/tint/757.wgsl:10:5 warning: use of deprecated language feature: [[offset]] has been replaced with [[size]] and [[align]]
-  [[offset(0)]] values : [[stride(4)]] array<f32>;
-    ^^^^^^
+kernel void tint_symbol(texture2d_array<float, access::sample> tint_symbol_2 [[texture(1)]], uint3 GlobalInvocationID [[thread_position_in_grid]], device Result& result [[buffer(3)]]) {
+  uint flatIndex = ((((2u * 2u) * GlobalInvocationID.z) + (2u * GlobalInvocationID.y)) + GlobalInvocationID.x);
+  flatIndex = (flatIndex * 1u);
+  float4 texel = tint_symbol_2.read(uint2(int2(GlobalInvocationID.xy)), 0, 0);
+  {
+    uint i = 0u;
+    {
+      bool tint_msl_is_first_1 = true;
+      for(;;) {
+        if (!tint_msl_is_first_1) {
+          i = (i + 1u);
+        }
+        tint_msl_is_first_1 = false;
 
+        if (!((i < 1u))) {
+          break;
+        }
+        result.values[(flatIndex + i)] = texel.r;
+      }
+    }
+  }
+  return;
+}
 
-
-Validation Failure:
-
-Compilation failed: 
-
-program_source:14:18: error: use of undeclared identifier 'myTexture'
-  float4 texel = myTexture.read(int2(GlobalInvocationID.xy), 0, 0);
-                 ^

test/bug/tint/827.wgsl.expected.msl

@@ -1,15 +1,13 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-bug/tint/827.wgsl:8:26 warning: use of deprecated language feature: declare access with var<storage, read_write> instead of using [[access]] decoration
-[[group(0), binding(1)]] var<storage> result : [[access(read_write)]] Result;
-                         ^^^
+using namespace metal;
+struct Result {
+  /* 0x0000 */ float values[1];
+};
 
+constant uint width = 128u;
+kernel void tint_symbol(depth2d<float, access::sample> tint_symbol_2 [[texture(0)]], uint3 GlobalInvocationId [[thread_position_in_grid]], device Result& result [[buffer(1)]]) {
+  result.values[((GlobalInvocationId.y * width) + GlobalInvocationId.x)] = tint_symbol_2.read(uint2(int2(int(GlobalInvocationId.x), int(GlobalInvocationId.y))), 0);
+  return;
+}
 
-
-Validation Failure:
-
-Compilation failed: 
-
-program_source:10:76: error: use of undeclared identifier 'tex'
-  result.values[((GlobalInvocationId.y * width) + GlobalInvocationId.x)] = tex.read(int2(int(GlobalInvocationId.x), int(GlobalInvocationId.y)), 0);
-                                                                           ^

test/bug/tint/870.spvasm.expected.hlsl

@@ -9,7 +9,7 @@ tint_array_wrapper tint_symbol_1(ByteAddressBuffer buffer, uint offset) {
 
 ByteAddressBuffer sspp962805860buildInformation : register(t2, space0);
 
-void main() {
+void main_1() {
   tint_array_wrapper orientation = {{0, 0, 0, 0, 0, 0}};
   const tint_array_wrapper x_23 = tint_symbol_1(sspp962805860buildInformation, 36u);
   orientation.arr[0] = x_23.arr[0u];
@@ -20,3 +20,8 @@ void main() {
   orientation.arr[5] = x_23.arr[5u];
   return;
 }
+
+void main() {
+  main_1();
+  return;
+}

test/bug/tint/870.spvasm.expected.msl

@@ -14,7 +14,7 @@ struct x_B4_BuildInformation {
   /* 0x0000 */ sspp962805860buildInformationS passthru;
 };
 
-fragment void tint_symbol(const device x_B4_BuildInformation& sspp962805860buildInformation [[buffer(2)]]) {
+void main_1(const device x_B4_BuildInformation& sspp962805860buildInformation) {
   tint_array_wrapper orientation = {};
   tint_array_wrapper const x_23 = sspp962805860buildInformation.passthru.orientation;
   orientation.arr[0] = x_23.arr[0u];
@@ -26,3 +26,8 @@ fragment void tint_symbol(const device x_B4_BuildInformation& sspp962805860build
   return;
 }
 
+fragment void tint_symbol(const device x_B4_BuildInformation& sspp962805860buildInformation [[buffer(2)]]) {
+  main_1(sspp962805860buildInformation);
+  return;
+}
+

test/bug/tint/870.spvasm.expected.spvasm

@@ -1,7 +1,7 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 46
+; Bound: 49
 ; Schema: 0
                OpCapability Shader
                OpMemoryModel Logical GLSL450
@@ -15,8 +15,9 @@
                OpMemberName %sspp962805860buildInformationS 2 "essence"
                OpMemberName %sspp962805860buildInformationS 3 "orientation"
                OpName %sspp962805860buildInformation "sspp962805860buildInformation"
-               OpName %main "main"
+               OpName %main_1 "main_1"
                OpName %orientation "orientation"
+               OpName %main "main"
                OpDecorate %x_B4_BuildInformation Block
                OpMemberDecorate %x_B4_BuildInformation 0 Offset 0
                OpMemberDecorate %sspp962805860buildInformationS 0 Offset 0
@@ -55,7 +56,7 @@
      %uint_4 = OpConstant %uint 4
       %int_5 = OpConstant %int 5
      %uint_5 = OpConstant %uint 5
-       %main = OpFunction %void None %11
+     %main_1 = OpFunction %void None %11
          %14 = OpLabel
 %orientation = OpVariable %_ptr_Function__arr_int_uint_6 Function %17
          %21 = OpAccessChain %_ptr_StorageBuffer__arr_int_uint_6 %sspp962805860buildInformation %uint_0 %uint_3
@@ -80,3 +81,8 @@
                OpStore %43 %45
                OpReturn
                OpFunctionEnd
+       %main = OpFunction %void None %11
+         %47 = OpLabel
+         %48 = OpFunctionCall %void %main_1
+               OpReturn
+               OpFunctionEnd

test/bug/tint/870.spvasm.expected.wgsl

@@ -14,8 +14,7 @@ struct x_B4_BuildInformation {
 
 [[group(0), binding(2)]] var<storage, read> sspp962805860buildInformation : x_B4_BuildInformation;
 
-[[stage(fragment)]]
-fn main() {
+fn main_1() {
   var orientation : array<i32, 6>;
   let x_23 : Arr = sspp962805860buildInformation.passthru.orientation;
   orientation[0] = x_23[0u];
@@ -26,3 +25,8 @@ fn main() {
   orientation[5] = x_23[5u];
   return;
 }
+
+[[stage(fragment)]]
+fn main() {
+  main_1();
+}

test/intrinsics/gen/textureSampleCompareLevel/011a8f.wgsl.expected.msl

@@ -1,14 +1,27 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
+using namespace metal;
+struct tint_symbol {
+  float4 value [[position]];
+};
 
+void textureSampleCompareLevel_011a8f(depth2d_array<float, access::sample> tint_symbol_2, sampler tint_symbol_3) {
+  float res = tint_symbol_2.sample_compare(tint_symbol_3, float2(), 1, 1.0f, level(0), int2());
+}
 
-Validation Failure:
+vertex tint_symbol vertex_main(depth2d_array<float, access::sample> tint_symbol_4 [[texture(0)]], sampler tint_symbol_5 [[sampler(1)]]) {
+  textureSampleCompareLevel_011a8f(tint_symbol_4, tint_symbol_5);
+  tint_symbol const tint_symbol_1 = {.value=float4()};
+  return tint_symbol_1;
+}
 
-Compilation failed: 
+fragment void fragment_main(depth2d_array<float, access::sample> tint_symbol_6 [[texture(0)]], sampler tint_symbol_7 [[sampler(1)]]) {
+  textureSampleCompareLevel_011a8f(tint_symbol_6, tint_symbol_7);
+  return;
+}
+
+kernel void compute_main(depth2d_array<float, access::sample> tint_symbol_8 [[texture(0)]], sampler tint_symbol_9 [[sampler(1)]]) {
+  textureSampleCompareLevel_011a8f(tint_symbol_8, tint_symbol_9);
+  return;
+}
 
-program_source:9:15: error: use of undeclared identifier 'arg_0'
-  float res = arg_0.sample_compare(arg_1, float2(), 1, 1.0f, level(0), int2());
-              ^
-program_source:9:36: error: use of undeclared identifier 'arg_1'
-  float res = arg_0.sample_compare(arg_1, float2(), 1, 1.0f, level(0), int2());
-                                   ^

test/intrinsics/gen/textureSampleCompareLevel/1116ed.wgsl.expected.msl

@@ -1,14 +1,27 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
+using namespace metal;
+struct tint_symbol {
+  float4 value [[position]];
+};
 
+void textureSampleCompareLevel_1116ed(depth2d_array<float, access::sample> tint_symbol_2, sampler tint_symbol_3) {
+  float res = tint_symbol_2.sample_compare(tint_symbol_3, float2(), 1, 1.0f, level(0));
+}
 
-Validation Failure:
+vertex tint_symbol vertex_main(depth2d_array<float, access::sample> tint_symbol_4 [[texture(0)]], sampler tint_symbol_5 [[sampler(1)]]) {
+  textureSampleCompareLevel_1116ed(tint_symbol_4, tint_symbol_5);
+  tint_symbol const tint_symbol_1 = {.value=float4()};
+  return tint_symbol_1;
+}
 
-Compilation failed: 
+fragment void fragment_main(depth2d_array<float, access::sample> tint_symbol_6 [[texture(0)]], sampler tint_symbol_7 [[sampler(1)]]) {
+  textureSampleCompareLevel_1116ed(tint_symbol_6, tint_symbol_7);
+  return;
+}
+
+kernel void compute_main(depth2d_array<float, access::sample> tint_symbol_8 [[texture(0)]], sampler tint_symbol_9 [[sampler(1)]]) {
+  textureSampleCompareLevel_1116ed(tint_symbol_8, tint_symbol_9);
+  return;
+}
 
-program_source:9:15: error: use of undeclared identifier 'arg_0'
-  float res = arg_0.sample_compare(arg_1, float2(), 1, 1.0f, level(0));
-              ^
-program_source:9:36: error: use of undeclared identifier 'arg_1'
-  float res = arg_0.sample_compare(arg_1, float2(), 1, 1.0f, level(0));
-                                   ^

test/intrinsics/gen/textureSampleCompareLevel/1568e3.wgsl.expected.msl

@@ -1,14 +1,27 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
+using namespace metal;
+struct tint_symbol {
+  float4 value [[position]];
+};
 
+void textureSampleCompareLevel_1568e3(depthcube<float, access::sample> tint_symbol_2, sampler tint_symbol_3) {
+  float res = tint_symbol_2.sample_compare(tint_symbol_3, float3(), 1.0f, level(0));
+}
 
-Validation Failure:
+vertex tint_symbol vertex_main(depthcube<float, access::sample> tint_symbol_4 [[texture(0)]], sampler tint_symbol_5 [[sampler(1)]]) {
+  textureSampleCompareLevel_1568e3(tint_symbol_4, tint_symbol_5);
+  tint_symbol const tint_symbol_1 = {.value=float4()};
+  return tint_symbol_1;
+}
 
-Compilation failed: 
+fragment void fragment_main(depthcube<float, access::sample> tint_symbol_6 [[texture(0)]], sampler tint_symbol_7 [[sampler(1)]]) {
+  textureSampleCompareLevel_1568e3(tint_symbol_6, tint_symbol_7);
+  return;
+}
+
+kernel void compute_main(depthcube<float, access::sample> tint_symbol_8 [[texture(0)]], sampler tint_symbol_9 [[sampler(1)]]) {
+  textureSampleCompareLevel_1568e3(tint_symbol_8, tint_symbol_9);
+  return;
+}
 
-program_source:9:15: error: use of undeclared identifier 'arg_0'
-  float res = arg_0.sample_compare(arg_1, float3(), 1.0f, level(0));
-              ^
-program_source:9:36: error: use of undeclared identifier 'arg_1'
-  float res = arg_0.sample_compare(arg_1, float3(), 1.0f, level(0));
-                                   ^

test/intrinsics/gen/textureSampleCompareLevel/2ad2b1.wgsl.expected.msl

@@ -1,14 +1,27 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
+using namespace metal;
+struct tint_symbol {
+  float4 value [[position]];
+};
 
+void textureSampleCompareLevel_2ad2b1(depth2d<float, access::sample> tint_symbol_2, sampler tint_symbol_3) {
+  float res = tint_symbol_2.sample_compare(tint_symbol_3, float2(), 1.0f, level(0));
+}
 
-Validation Failure:
+vertex tint_symbol vertex_main(depth2d<float, access::sample> tint_symbol_4 [[texture(0)]], sampler tint_symbol_5 [[sampler(1)]]) {
+  textureSampleCompareLevel_2ad2b1(tint_symbol_4, tint_symbol_5);
+  tint_symbol const tint_symbol_1 = {.value=float4()};
+  return tint_symbol_1;
+}
 
-Compilation failed: 
+fragment void fragment_main(depth2d<float, access::sample> tint_symbol_6 [[texture(0)]], sampler tint_symbol_7 [[sampler(1)]]) {
+  textureSampleCompareLevel_2ad2b1(tint_symbol_6, tint_symbol_7);
+  return;
+}
+
+kernel void compute_main(depth2d<float, access::sample> tint_symbol_8 [[texture(0)]], sampler tint_symbol_9 [[sampler(1)]]) {
+  textureSampleCompareLevel_2ad2b1(tint_symbol_8, tint_symbol_9);
+  return;
+}
 
-program_source:9:15: error: use of undeclared identifier 'arg_0'
-  float res = arg_0.sample_compare(arg_1, float2(), 1.0f, level(0));
-              ^
-program_source:9:36: error: use of undeclared identifier 'arg_1'
-  float res = arg_0.sample_compare(arg_1, float2(), 1.0f, level(0));
-                                   ^

test/intrinsics/gen/textureSampleCompareLevel/4cf3a2.wgsl.expected.msl

@@ -1,14 +1,27 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
+using namespace metal;
+struct tint_symbol {
+  float4 value [[position]];
+};
 
+void textureSampleCompareLevel_4cf3a2(depthcube_array<float, access::sample> tint_symbol_2, sampler tint_symbol_3) {
+  float res = tint_symbol_2.sample_compare(tint_symbol_3, float3(), 1, 1.0f, level(0));
+}
 
-Validation Failure:
+vertex tint_symbol vertex_main(depthcube_array<float, access::sample> tint_symbol_4 [[texture(0)]], sampler tint_symbol_5 [[sampler(1)]]) {
+  textureSampleCompareLevel_4cf3a2(tint_symbol_4, tint_symbol_5);
+  tint_symbol const tint_symbol_1 = {.value=float4()};
+  return tint_symbol_1;
+}
 
-Compilation failed: 
+fragment void fragment_main(depthcube_array<float, access::sample> tint_symbol_6 [[texture(0)]], sampler tint_symbol_7 [[sampler(1)]]) {
+  textureSampleCompareLevel_4cf3a2(tint_symbol_6, tint_symbol_7);
+  return;
+}
+
+kernel void compute_main(depthcube_array<float, access::sample> tint_symbol_8 [[texture(0)]], sampler tint_symbol_9 [[sampler(1)]]) {
+  textureSampleCompareLevel_4cf3a2(tint_symbol_8, tint_symbol_9);
+  return;
+}
 
-program_source:9:15: error: use of undeclared identifier 'arg_0'
-  float res = arg_0.sample_compare(arg_1, float3(), 1, 1.0f, level(0));
-              ^
-program_source:9:36: error: use of undeclared identifier 'arg_1'
-  float res = arg_0.sample_compare(arg_1, float3(), 1, 1.0f, level(0));
-                                   ^

test/intrinsics/gen/textureSampleCompareLevel/f8121c.wgsl.expected.msl

@@ -1,14 +1,27 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
+using namespace metal;
+struct tint_symbol {
+  float4 value [[position]];
+};
 
+void textureSampleCompareLevel_f8121c(depth2d<float, access::sample> tint_symbol_2, sampler tint_symbol_3) {
+  float res = tint_symbol_2.sample_compare(tint_symbol_3, float2(), 1.0f, level(0), int2());
+}
 
-Validation Failure:
+vertex tint_symbol vertex_main(depth2d<float, access::sample> tint_symbol_4 [[texture(0)]], sampler tint_symbol_5 [[sampler(1)]]) {
+  textureSampleCompareLevel_f8121c(tint_symbol_4, tint_symbol_5);
+  tint_symbol const tint_symbol_1 = {.value=float4()};
+  return tint_symbol_1;
+}
 
-Compilation failed: 
+fragment void fragment_main(depth2d<float, access::sample> tint_symbol_6 [[texture(0)]], sampler tint_symbol_7 [[sampler(1)]]) {
+  textureSampleCompareLevel_f8121c(tint_symbol_6, tint_symbol_7);
+  return;
+}
+
+kernel void compute_main(depth2d<float, access::sample> tint_symbol_8 [[texture(0)]], sampler tint_symbol_9 [[sampler(1)]]) {
+  textureSampleCompareLevel_f8121c(tint_symbol_8, tint_symbol_9);
+  return;
+}
 
-program_source:9:15: error: use of undeclared identifier 'arg_0'
-  float res = arg_0.sample_compare(arg_1, float2(), 1.0f, level(0), int2());
-              ^
-program_source:9:36: error: use of undeclared identifier 'arg_1'
-  float res = arg_0.sample_compare(arg_1, float2(), 1.0f, level(0), int2());
-                                   ^

test/ptr_ref/access/matrix.spvasm.expected.hlsl

@@ -1,7 +1,12 @@
-[numthreads(1, 1, 1)]
-void main() {
+void main_1() {
   float3x3 m = float3x3(float3(0.0f, 0.0f, 0.0f), float3(0.0f, 0.0f, 0.0f), float3(0.0f, 0.0f, 0.0f));
   m = float3x3(float3(1.0f, 2.0f, 3.0f), float3(4.0f, 5.0f, 6.0f), float3(7.0f, 8.0f, 9.0f));
   m[1] = float3(5.0f, 5.0f, 5.0f);
   return;
 }
+
+[numthreads(1, 1, 1)]
+void main() {
+  main_1();
+  return;
+}

test/ptr_ref/access/matrix.spvasm.expected.msl

@@ -1,10 +1,15 @@
 #include <metal_stdlib>
 
 using namespace metal;
-kernel void tint_symbol() {
+void main_1() {
   float3x3 m = float3x3(float3(0.0f, 0.0f, 0.0f), float3(0.0f, 0.0f, 0.0f), float3(0.0f, 0.0f, 0.0f));
   m = float3x3(float3(1.0f, 2.0f, 3.0f), float3(4.0f, 5.0f, 6.0f), float3(7.0f, 8.0f, 9.0f));
   m[1] = float3(5.0f, 5.0f, 5.0f);
   return;
 }
 
+kernel void tint_symbol() {
+  main_1();
+  return;
+}
+

test/ptr_ref/access/matrix.spvasm.expected.spvasm

@@ -1,14 +1,15 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 32
+; Bound: 35
 ; Schema: 0
                OpCapability Shader
                OpMemoryModel Logical GLSL450
                OpEntryPoint GLCompute %main "main"
                OpExecutionMode %main LocalSize 1 1 1
-               OpName %main "main"
+               OpName %main_1 "main_1"
                OpName %m "m"
+               OpName %main "main"
        %void = OpTypeVoid
           %1 = OpTypeFunction %void
       %float = OpTypeFloat 32
@@ -36,7 +37,7 @@
       %int_1 = OpConstant %int 1
 %_ptr_Function_v3float = OpTypePointer Function %v3float
          %31 = OpConstantComposite %v3float %float_5 %float_5 %float_5
-       %main = OpFunction %void None %1
+     %main_1 = OpFunction %void None %1
           %4 = OpLabel
           %m = OpVariable %_ptr_Function_mat3v3float Function %13
                OpStore %m %10
@@ -45,3 +46,8 @@
                OpStore %30 %31
                OpReturn
                OpFunctionEnd
+       %main = OpFunction %void None %1
+         %33 = OpLabel
+         %34 = OpFunctionCall %void %main_1
+               OpReturn
+               OpFunctionEnd

test/ptr_ref/access/matrix.spvasm.expected.wgsl

@@ -1,7 +1,11 @@
-[[stage(compute)]]
-fn main() {
+fn main_1() {
   var m : mat3x3<f32> = mat3x3<f32>(vec3<f32>(0.0, 0.0, 0.0), vec3<f32>(0.0, 0.0, 0.0), vec3<f32>(0.0, 0.0, 0.0));
   m = mat3x3<f32>(vec3<f32>(1.0, 2.0, 3.0), vec3<f32>(4.0, 5.0, 6.0), vec3<f32>(7.0, 8.0, 9.0));
   m[1] = vec3<f32>(5.0, 5.0, 5.0);
   return;
 }
+
+[[stage(compute)]]
+fn main() {
+  main_1();
+}

test/ptr_ref/access/vector.spvasm.expected.hlsl

@@ -1,7 +1,12 @@
-[numthreads(1, 1, 1)]
-void main() {
+void main_1() {
   float3 v = float3(0.0f, 0.0f, 0.0f);
   v = float3(1.0f, 2.0f, 3.0f);
   v.y = 5.0f;
   return;
 }
+
+[numthreads(1, 1, 1)]
+void main() {
+  main_1();
+  return;
+}

test/ptr_ref/access/vector.spvasm.expected.msl

@@ -1,10 +1,15 @@
 #include <metal_stdlib>
 
 using namespace metal;
-kernel void tint_symbol() {
+void main_1() {
   float3 v = float3(0.0f, 0.0f, 0.0f);
   v = float3(1.0f, 2.0f, 3.0f);
   v.y = 5.0f;
   return;
 }
 
+kernel void tint_symbol() {
+  main_1();
+  return;
+}
+

test/ptr_ref/access/vector.spvasm.expected.spvasm

@@ -1,14 +1,15 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 21
+; Bound: 24
 ; Schema: 0
                OpCapability Shader
                OpMemoryModel Logical GLSL450
                OpEntryPoint GLCompute %main "main"
                OpExecutionMode %main LocalSize 1 1 1
-               OpName %main "main"
+               OpName %main_1 "main_1"
                OpName %v "v"
+               OpName %main "main"
        %void = OpTypeVoid
           %1 = OpTypeFunction %void
       %float = OpTypeFloat 32
@@ -25,7 +26,7 @@
      %uint_1 = OpConstant %uint 1
 %_ptr_Function_float = OpTypePointer Function %float
     %float_5 = OpConstant %float 5
-       %main = OpFunction %void None %1
+     %main_1 = OpFunction %void None %1
           %4 = OpLabel
           %v = OpVariable %_ptr_Function_v3float Function %11
                OpStore %v %8
@@ -34,3 +35,8 @@
                OpStore %19 %float_5
                OpReturn
                OpFunctionEnd
+       %main = OpFunction %void None %1
+         %22 = OpLabel
+         %23 = OpFunctionCall %void %main_1
+               OpReturn
+               OpFunctionEnd

test/ptr_ref/access/vector.spvasm.expected.wgsl

@@ -1,7 +1,11 @@
-[[stage(compute)]]
-fn main() {
+fn main_1() {
   var v : vec3<f32> = vec3<f32>(0.0, 0.0, 0.0);
   v = vec3<f32>(1.0, 2.0, 3.0);
   v.y = 5.0;
   return;
 }
+
+[[stage(compute)]]
+fn main() {
+  main_1();
+}

test/ptr_ref/copy/ptr_copy.spvasm.expected.hlsl

@@ -1,5 +1,10 @@
-[numthreads(1, 1, 1)]
-void main() {
+void main_1() {
   uint x_10 = 0u;
   return;
 }
+
+[numthreads(1, 1, 1)]
+void main() {
+  main_1();
+  return;
+}

test/ptr_ref/copy/ptr_copy.spvasm.expected.msl

@@ -1,10 +1,15 @@
 #include <metal_stdlib>
 
 using namespace metal;
-kernel void tint_symbol() {
+void main_1() {
   uint x_10 = 0u;
   thread uint* const x_1 = &(x_10);
   thread uint* const x_2 = x_1;
   return;
 }
 
+kernel void tint_symbol() {
+  main_1();
+  return;
+}
+

test/ptr_ref/copy/ptr_copy.spvasm.expected.spvasm

@@ -1,21 +1,27 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 9
+; Bound: 12
 ; Schema: 0
                OpCapability Shader
                OpMemoryModel Logical GLSL450
                OpEntryPoint GLCompute %main "main"
                OpExecutionMode %main LocalSize 1 1 1
-               OpName %main "main"
+               OpName %main_1 "main_1"
                OpName %x_10 "x_10"
+               OpName %main "main"
        %void = OpTypeVoid
           %1 = OpTypeFunction %void
        %uint = OpTypeInt 32 0
 %_ptr_Function_uint = OpTypePointer Function %uint
           %8 = OpConstantNull %uint
-       %main = OpFunction %void None %1
+     %main_1 = OpFunction %void None %1
           %4 = OpLabel
        %x_10 = OpVariable %_ptr_Function_uint Function %8
                OpReturn
                OpFunctionEnd
+       %main = OpFunction %void None %1
+         %10 = OpLabel
+         %11 = OpFunctionCall %void %main_1
+               OpReturn
+               OpFunctionEnd

test/ptr_ref/copy/ptr_copy.spvasm.expected.wgsl

@@ -1,7 +1,11 @@
-[[stage(compute)]]
-fn main() {
+fn main_1() {
   var x_10 : u32;
   let x_1 : ptr<function, u32> = &(x_10);
   let x_2 : ptr<function, u32> = x_1;
   return;
 }
+
+[[stage(compute)]]
+fn main() {
+  main_1();
+}

test/ptr_ref/load/global/i32.spvasm.expected.hlsl

@@ -1,8 +1,13 @@
 static int I = 0;
 
-[numthreads(1, 1, 1)]
-void main() {
+void main_1() {
   const int x_9 = I;
   const int x_11 = (x_9 + 1);
   return;
 }
+
+[numthreads(1, 1, 1)]
+void main() {
+  main_1();
+  return;
+}

test/ptr_ref/load/global/i32.spvasm.expected.msl

@@ -1,10 +1,15 @@
 #include <metal_stdlib>
 
 using namespace metal;
-kernel void tint_symbol() {
-  thread int tint_symbol_1 = 0;
-  int const x_9 = tint_symbol_1;
+void main_1(thread int* const tint_symbol_1) {
+  int const x_9 = *(tint_symbol_1);
   int const x_11 = (x_9 + 1);
   return;
 }
 
+kernel void tint_symbol() {
+  thread int tint_symbol_2 = 0;
+  main_1(&(tint_symbol_2));
+  return;
+}
+

test/ptr_ref/load/global/i32.spvasm.expected.spvasm

@@ -1,13 +1,14 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 12
+; Bound: 15
 ; Schema: 0
                OpCapability Shader
                OpMemoryModel Logical GLSL450
                OpEntryPoint GLCompute %main "main"
                OpExecutionMode %main LocalSize 1 1 1
                OpName %I "I"
+               OpName %main_1 "main_1"
                OpName %main "main"
         %int = OpTypeInt 32 1
       %int_0 = OpConstant %int 0
@@ -16,9 +17,14 @@
        %void = OpTypeVoid
           %5 = OpTypeFunction %void
       %int_1 = OpConstant %int 1
-       %main = OpFunction %void None %5
+     %main_1 = OpFunction %void None %5
           %8 = OpLabel
           %9 = OpLoad %int %I
          %11 = OpIAdd %int %9 %int_1
                OpReturn
                OpFunctionEnd
+       %main = OpFunction %void None %5
+         %13 = OpLabel
+         %14 = OpFunctionCall %void %main_1
+               OpReturn
+               OpFunctionEnd

test/ptr_ref/load/global/i32.spvasm.expected.wgsl

@@ -1,8 +1,12 @@
 var<private> I : i32 = 0;
 
-[[stage(compute)]]
-fn main() {
+fn main_1() {
   let x_9 : i32 = I;
   let x_11 : i32 = (x_9 + 1);
   return;
 }
+
+[[stage(compute)]]
+fn main() {
+  main_1();
+}

test/ptr_ref/load/global/struct_field.spvasm.expected.hlsl

@@ -4,10 +4,15 @@ struct S {
 
 static S V;
 
-[numthreads(1, 1, 1)]
-void main() {
+void main_1() {
   int i = 0;
   const int x_15 = V.i;
   i = x_15;
   return;
 }
+
+[numthreads(1, 1, 1)]
+void main() {
+  main_1();
+  return;
+}

test/ptr_ref/load/global/struct_field.spvasm.expected.msl

@@ -5,11 +5,16 @@ struct S {
   int i;
 };
 
-kernel void tint_symbol() {
-  thread S tint_symbol_1 = {};
+void main_1(thread S* const tint_symbol_1) {
   int i = 0;
-  int const x_15 = tint_symbol_1.i;
+  int const x_15 = (*(tint_symbol_1)).i;
   i = x_15;
   return;
 }
 
+kernel void tint_symbol() {
+  thread S tint_symbol_2 = {};
+  main_1(&(tint_symbol_2));
+  return;
+}
+

test/ptr_ref/load/global/struct_field.spvasm.expected.spvasm

@@ -1,7 +1,7 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 18
+; Bound: 21
 ; Schema: 0
                OpCapability Shader
                OpMemoryModel Logical GLSL450
@@ -10,8 +10,9 @@
                OpName %S "S"
                OpMemberName %S 0 "i"
                OpName %V "V"
-               OpName %main "main"
+               OpName %main_1 "main_1"
                OpName %i "i"
+               OpName %main "main"
                OpMemberDecorate %S 0 Offset 0
         %int = OpTypeInt 32 1
           %S = OpTypeStruct %int
@@ -25,7 +26,7 @@
        %uint = OpTypeInt 32 0
      %uint_0 = OpConstant %uint 0
 %_ptr_Private_int = OpTypePointer Private %int
-       %main = OpFunction %void None %6
+     %main_1 = OpFunction %void None %6
           %9 = OpLabel
           %i = OpVariable %_ptr_Function_int Function %12
          %16 = OpAccessChain %_ptr_Private_int %V %uint_0
@@ -33,3 +34,8 @@
                OpStore %i %17
                OpReturn
                OpFunctionEnd
+       %main = OpFunction %void None %6
+         %19 = OpLabel
+         %20 = OpFunctionCall %void %main_1
+               OpReturn
+               OpFunctionEnd

test/ptr_ref/load/global/struct_field.spvasm.expected.wgsl

@@ -4,10 +4,14 @@ struct S {
 
 var<private> V : S;
 
-[[stage(compute)]]
-fn main() {
+fn main_1() {
   var i : i32;
   let x_15 : i32 = V.i;
   i = x_15;
   return;
 }
+
+[[stage(compute)]]
+fn main() {
+  main_1();
+}

test/ptr_ref/load/local/i32.spvasm.expected.hlsl

@@ -1,8 +1,13 @@
-[numthreads(1, 1, 1)]
-void main() {
+void main_1() {
   int i = 0;
   i = 123;
   const int x_10 = i;
   const int x_12 = (x_10 + 1);
   return;
 }
+
+[numthreads(1, 1, 1)]
+void main() {
+  main_1();
+  return;
+}

test/ptr_ref/load/local/i32.spvasm.expected.msl

@@ -1,7 +1,7 @@
 #include <metal_stdlib>
 
 using namespace metal;
-kernel void tint_symbol() {
+void main_1() {
   int i = 0;
   i = 123;
   int const x_10 = i;
@@ -9,3 +9,8 @@ kernel void tint_symbol() {
   return;
 }
 
+kernel void tint_symbol() {
+  main_1();
+  return;
+}
+

test/ptr_ref/load/local/i32.spvasm.expected.spvasm

@@ -1,14 +1,15 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 14
+; Bound: 17
 ; Schema: 0
                OpCapability Shader
                OpMemoryModel Logical GLSL450
                OpEntryPoint GLCompute %main "main"
                OpExecutionMode %main LocalSize 1 1 1
-               OpName %main "main"
+               OpName %main_1 "main_1"
                OpName %i "i"
+               OpName %main "main"
        %void = OpTypeVoid
           %1 = OpTypeFunction %void
         %int = OpTypeInt 32 1
@@ -17,7 +18,7 @@
           %9 = OpConstantNull %int
     %int_123 = OpConstant %int 123
       %int_1 = OpConstant %int 1
-       %main = OpFunction %void None %1
+     %main_1 = OpFunction %void None %1
           %4 = OpLabel
           %i = OpVariable %_ptr_Function_int Function %9
                OpStore %i %int_0
@@ -26,3 +27,8 @@
          %13 = OpIAdd %int %11 %int_1
                OpReturn
                OpFunctionEnd
+       %main = OpFunction %void None %1
+         %15 = OpLabel
+         %16 = OpFunctionCall %void %main_1
+               OpReturn
+               OpFunctionEnd

test/ptr_ref/load/local/i32.spvasm.expected.wgsl

@@ -1,8 +1,12 @@
-[[stage(compute)]]
-fn main() {
+fn main_1() {
   var i : i32 = 0;
   i = 123;
   let x_10 : i32 = i;
   let x_12 : i32 = (x_10 + 1);
   return;
 }
+
+[[stage(compute)]]
+fn main() {
+  main_1();
+}

test/ptr_ref/load/local/struct_field.spvasm.expected.hlsl

@@ -2,11 +2,16 @@ struct S {
   int i;
 };
 
-[numthreads(1, 1, 1)]
-void main() {
+void main_1() {
   int i = 0;
   S V = {0};
   const int x_14 = V.i;
   i = x_14;
   return;
 }
+
+[numthreads(1, 1, 1)]
+void main() {
+  main_1();
+  return;
+}

test/ptr_ref/load/local/struct_field.spvasm.expected.msl

@@ -5,7 +5,7 @@ struct S {
   int i;
 };
 
-kernel void tint_symbol() {
+void main_1() {
   int i = 0;
   S V = {};
   int const x_14 = V.i;
@@ -13,3 +13,8 @@ kernel void tint_symbol() {
   return;
 }
 
+kernel void tint_symbol() {
+  main_1();
+  return;
+}
+

test/ptr_ref/load/local/struct_field.spvasm.expected.spvasm

@@ -1,17 +1,18 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 17
+; Bound: 20
 ; Schema: 0
                OpCapability Shader
                OpMemoryModel Logical GLSL450
                OpEntryPoint GLCompute %main "main"
                OpExecutionMode %main LocalSize 1 1 1
-               OpName %main "main"
+               OpName %main_1 "main_1"
                OpName %i "i"
                OpName %S "S"
                OpMemberName %S 0 "i"
                OpName %V "V"
+               OpName %main "main"
                OpMemberDecorate %S 0 Offset 0
        %void = OpTypeVoid
           %1 = OpTypeFunction %void
@@ -23,7 +24,7 @@
          %12 = OpConstantNull %S
        %uint = OpTypeInt 32 0
      %uint_0 = OpConstant %uint 0
-       %main = OpFunction %void None %1
+     %main_1 = OpFunction %void None %1
           %4 = OpLabel
           %i = OpVariable %_ptr_Function_int Function %8
           %V = OpVariable %_ptr_Function_S Function %12
@@ -32,3 +33,8 @@
                OpStore %i %16
                OpReturn
                OpFunctionEnd
+       %main = OpFunction %void None %1
+         %18 = OpLabel
+         %19 = OpFunctionCall %void %main_1
+               OpReturn
+               OpFunctionEnd

test/ptr_ref/load/local/struct_field.spvasm.expected.wgsl

@@ -2,11 +2,15 @@ struct S {
   i : i32;
 };
 
-[[stage(compute)]]
-fn main() {
+fn main_1() {
   var i : i32;
   var V : S;
   let x_14 : i32 = V.i;
   i = x_14;
   return;
 }
+
+[[stage(compute)]]
+fn main() {
+  main_1();
+}

test/ptr_ref/load/param/ptr.spvasm.expected.hlsl

@@ -3,11 +3,16 @@ int func(int value, inout int pointer) {
   return (value + x_9);
 }
 
-[numthreads(1, 1, 1)]
-void main() {
+void main_1() {
   int i = 0;
   i = 123;
   const int x_19 = i;
   const int x_18 = func(x_19, i);
   return;
 }
+
+[numthreads(1, 1, 1)]
+void main() {
+  main_1();
+  return;
+}

test/ptr_ref/load/param/ptr.spvasm.expected.msl

@@ -6,7 +6,7 @@ int func(int value, thread int* const pointer) {
   return (value + x_9);
 }
 
-kernel void tint_symbol() {
+void main_1() {
   int i = 0;
   i = 123;
   int const x_19 = i;
@@ -14,3 +14,8 @@ kernel void tint_symbol() {
   return;
 }
 
+kernel void tint_symbol() {
+  main_1();
+  return;
+}
+

test/ptr_ref/load/param/ptr.spvasm.expected.spvasm

@@ -1,7 +1,7 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 22
+; Bound: 25
 ; Schema: 0
                OpCapability Shader
                OpMemoryModel Logical GLSL450
@@ -10,8 +10,9 @@
                OpName %func "func"
                OpName %value "value"
                OpName %pointer "pointer"
-               OpName %main "main"
+               OpName %main_1 "main_1"
                OpName %i "i"
+               OpName %main "main"
         %int = OpTypeInt 32 1
 %_ptr_Function_int = OpTypePointer Function %int
           %1 = OpTypeFunction %int %int %_ptr_Function_int
@@ -28,7 +29,7 @@
          %10 = OpIAdd %int %value %9
                OpReturnValue %10
                OpFunctionEnd
-       %main = OpFunction %void None %11
+     %main_1 = OpFunction %void None %11
          %14 = OpLabel
           %i = OpVariable %_ptr_Function_int Function %17
                OpStore %i %int_0
@@ -37,3 +38,8 @@
          %20 = OpFunctionCall %int %func %19 %i
                OpReturn
                OpFunctionEnd
+       %main = OpFunction %void None %11
+         %23 = OpLabel
+         %24 = OpFunctionCall %void %main_1
+               OpReturn
+               OpFunctionEnd

test/ptr_ref/load/param/ptr.spvasm.expected.wgsl

@@ -3,11 +3,15 @@ fn func(value : i32, pointer : ptr<function, i32>) -> i32 {
   return (value + x_9);
 }
 
-[[stage(compute)]]
-fn main() {
+fn main_1() {
   var i : i32 = 0;
   i = 123;
   let x_19 : i32 = i;
   let x_18 : i32 = func(x_19, &(i));
   return;
 }
+
+[[stage(compute)]]
+fn main() {
+  main_1();
+}

test/ptr_ref/store/global/i32.spvasm.expected.hlsl

@@ -1,8 +1,13 @@
 static int I = 0;
 
-[numthreads(1, 1, 1)]
-void main() {
+void main_1() {
   I = 123;
   I = ((100 + 20) + 3);
   return;
 }
+
+[numthreads(1, 1, 1)]
+void main() {
+  main_1();
+  return;
+}

test/ptr_ref/store/global/i32.spvasm.expected.msl

@@ -1,10 +1,15 @@
 #include <metal_stdlib>
 
 using namespace metal;
-kernel void tint_symbol() {
-  thread int tint_symbol_1 = 0;
-  tint_symbol_1 = 123;
-  tint_symbol_1 = ((100 + 20) + 3);
+void main_1(thread int* const tint_symbol_1) {
+  *(tint_symbol_1) = 123;
+  *(tint_symbol_1) = ((100 + 20) + 3);
+  return;
+}
+
+kernel void tint_symbol() {
+  thread int tint_symbol_2 = 0;
+  main_1(&(tint_symbol_2));
   return;
 }
 

test/ptr_ref/store/global/i32.spvasm.expected.spvasm

@@ -1,13 +1,14 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 15
+; Bound: 18
 ; Schema: 0
                OpCapability Shader
                OpMemoryModel Logical GLSL450
                OpEntryPoint GLCompute %main "main"
                OpExecutionMode %main LocalSize 1 1 1
                OpName %I "I"
+               OpName %main_1 "main_1"
                OpName %main "main"
         %int = OpTypeInt 32 1
       %int_0 = OpConstant %int 0
@@ -19,7 +20,7 @@
     %int_100 = OpConstant %int 100
      %int_20 = OpConstant %int 20
       %int_3 = OpConstant %int 3
-       %main = OpFunction %void None %5
+     %main_1 = OpFunction %void None %5
           %8 = OpLabel
                OpStore %I %int_123
          %12 = OpIAdd %int %int_100 %int_20
@@ -27,3 +28,8 @@
                OpStore %I %14
                OpReturn
                OpFunctionEnd
+       %main = OpFunction %void None %5
+         %16 = OpLabel
+         %17 = OpFunctionCall %void %main_1
+               OpReturn
+               OpFunctionEnd

test/ptr_ref/store/global/i32.spvasm.expected.wgsl

@@ -1,8 +1,12 @@
 var<private> I : i32 = 0;
 
-[[stage(compute)]]
-fn main() {
+fn main_1() {
   I = 123;
   I = ((100 + 20) + 3);
   return;
 }
+
+[[stage(compute)]]
+fn main() {
+  main_1();
+}

test/ptr_ref/store/global/struct_field.spvasm.expected.hlsl

@@ -4,8 +4,13 @@ struct S {
 
 static S V;
 
-[numthreads(1, 1, 1)]
-void main() {
+void main_1() {
   V.i = 5;
   return;
 }
+
+[numthreads(1, 1, 1)]
+void main() {
+  main_1();
+  return;
+}

test/ptr_ref/store/global/struct_field.spvasm.expected.msl

@@ -5,9 +5,14 @@ struct S {
   int i;
 };
 
-kernel void tint_symbol() {
-  thread S tint_symbol_1 = {};
-  tint_symbol_1.i = 5;
+void main_1(thread S* const tint_symbol_1) {
+  (*(tint_symbol_1)).i = 5;
+  return;
+}
+
+kernel void tint_symbol() {
+  thread S tint_symbol_2 = {};
+  main_1(&(tint_symbol_2));
   return;
 }
 

test/ptr_ref/store/global/struct_field.spvasm.expected.spvasm

@@ -1,7 +1,7 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 15
+; Bound: 18
 ; Schema: 0
                OpCapability Shader
                OpMemoryModel Logical GLSL450
@@ -10,6 +10,7 @@
                OpName %S "S"
                OpMemberName %S 0 "i"
                OpName %V "V"
+               OpName %main_1 "main_1"
                OpName %main "main"
                OpMemberDecorate %S 0 Offset 0
         %int = OpTypeInt 32 1
@@ -23,9 +24,14 @@
      %uint_0 = OpConstant %uint 0
 %_ptr_Private_int = OpTypePointer Private %int
       %int_5 = OpConstant %int 5
-       %main = OpFunction %void None %6
+     %main_1 = OpFunction %void None %6
           %9 = OpLabel
          %13 = OpAccessChain %_ptr_Private_int %V %uint_0
                OpStore %13 %int_5
                OpReturn
                OpFunctionEnd
+       %main = OpFunction %void None %6
+         %16 = OpLabel
+         %17 = OpFunctionCall %void %main_1
+               OpReturn
+               OpFunctionEnd

test/ptr_ref/store/global/struct_field.spvasm.expected.wgsl

@@ -4,8 +4,12 @@ struct S {
 
 var<private> V : S;
 
-[[stage(compute)]]
-fn main() {
+fn main_1() {
   V.i = 5;
   return;
 }
+
+[[stage(compute)]]
+fn main() {
+  main_1();
+}

test/ptr_ref/store/local/i32.spvasm.expected.hlsl

@@ -1,8 +1,13 @@
-[numthreads(1, 1, 1)]
-void main() {
+void main_1() {
   int i = 0;
   i = 123;
   i = 123;
   i = ((100 + 20) + 3);
   return;
 }
+
+[numthreads(1, 1, 1)]
+void main() {
+  main_1();
+  return;
+}

test/ptr_ref/store/local/i32.spvasm.expected.msl

@@ -1,7 +1,7 @@
 #include <metal_stdlib>
 
 using namespace metal;
-kernel void tint_symbol() {
+void main_1() {
   int i = 0;
   i = 123;
   i = 123;
@@ -9,3 +9,8 @@ kernel void tint_symbol() {
   return;
 }
 
+kernel void tint_symbol() {
+  main_1();
+  return;
+}
+

test/ptr_ref/store/local/i32.spvasm.expected.spvasm

@@ -1,14 +1,15 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 16
+; Bound: 19
 ; Schema: 0
                OpCapability Shader
                OpMemoryModel Logical GLSL450
                OpEntryPoint GLCompute %main "main"
                OpExecutionMode %main LocalSize 1 1 1
-               OpName %main "main"
+               OpName %main_1 "main_1"
                OpName %i "i"
+               OpName %main "main"
        %void = OpTypeVoid
           %1 = OpTypeFunction %void
         %int = OpTypeInt 32 1
@@ -19,7 +20,7 @@
     %int_100 = OpConstant %int 100
      %int_20 = OpConstant %int 20
       %int_3 = OpConstant %int 3
-       %main = OpFunction %void None %1
+     %main_1 = OpFunction %void None %1
           %4 = OpLabel
           %i = OpVariable %_ptr_Function_int Function %9
                OpStore %i %int_0
@@ -30,3 +31,8 @@
                OpStore %i %15
                OpReturn
                OpFunctionEnd
+       %main = OpFunction %void None %1
+         %17 = OpLabel
+         %18 = OpFunctionCall %void %main_1
+               OpReturn
+               OpFunctionEnd

test/ptr_ref/store/local/i32.spvasm.expected.wgsl

@@ -1,8 +1,12 @@
-[[stage(compute)]]
-fn main() {
+fn main_1() {
   var i : i32 = 0;
   i = 123;
   i = 123;
   i = ((100 + 20) + 3);
   return;
 }
+
+[[stage(compute)]]
+fn main() {
+  main_1();
+}

test/ptr_ref/store/local/struct_field.spvasm.expected.hlsl

@@ -2,9 +2,14 @@ struct S {
   int i;
 };
 
-[numthreads(1, 1, 1)]
-void main() {
+void main_1() {
   S V = {0};
   V.i = 5;
   return;
 }
+
+[numthreads(1, 1, 1)]
+void main() {
+  main_1();
+  return;
+}

test/ptr_ref/store/local/struct_field.spvasm.expected.msl

@@ -5,9 +5,14 @@ struct S {
   int i;
 };
 
-kernel void tint_symbol() {
+void main_1() {
   S V = {};
   V.i = 5;
   return;
 }
 
+kernel void tint_symbol() {
+  main_1();
+  return;
+}
+

test/ptr_ref/store/local/struct_field.spvasm.expected.spvasm

@@ -1,16 +1,17 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 15
+; Bound: 18
 ; Schema: 0
                OpCapability Shader
                OpMemoryModel Logical GLSL450
                OpEntryPoint GLCompute %main "main"
                OpExecutionMode %main LocalSize 1 1 1
-               OpName %main "main"
+               OpName %main_1 "main_1"
                OpName %S "S"
                OpMemberName %S 0 "i"
                OpName %V "V"
+               OpName %main "main"
                OpMemberDecorate %S 0 Offset 0
        %void = OpTypeVoid
           %1 = OpTypeFunction %void
@@ -22,10 +23,15 @@
      %uint_0 = OpConstant %uint 0
 %_ptr_Function_int = OpTypePointer Function %int
       %int_5 = OpConstant %int 5
-       %main = OpFunction %void None %1
+     %main_1 = OpFunction %void None %1
           %4 = OpLabel
           %V = OpVariable %_ptr_Function_S Function %9
          %13 = OpAccessChain %_ptr_Function_int %V %uint_0
                OpStore %13 %int_5
                OpReturn
                OpFunctionEnd
+       %main = OpFunction %void None %1
+         %16 = OpLabel
+         %17 = OpFunctionCall %void %main_1
+               OpReturn
+               OpFunctionEnd

test/ptr_ref/store/local/struct_field.spvasm.expected.wgsl

@@ -2,9 +2,13 @@ struct S {
   i : i32;
 };
 
-[[stage(compute)]]
-fn main() {
+fn main_1() {
   var V : S;
   V.i = 5;
   return;
 }
+
+[[stage(compute)]]
+fn main() {
+  main_1();
+}

test/ptr_ref/store/param/ptr.spvasm.expected.hlsl

@@ -3,10 +3,15 @@ void func(int value, inout int pointer) {
   return;
 }
 
-[numthreads(1, 1, 1)]
-void main() {
+void main_1() {
   int i = 0;
   i = 123;
   func(123, i);
   return;
 }
+
+[numthreads(1, 1, 1)]
+void main() {
+  main_1();
+  return;
+}

test/ptr_ref/store/param/ptr.spvasm.expected.msl

@@ -6,10 +6,15 @@ void func(int value, thread int* const pointer) {
   return;
 }
 
-kernel void tint_symbol() {
+void main_1() {
   int i = 0;
   i = 123;
   func(123, &(i));
   return;
 }
 
+kernel void tint_symbol() {
+  main_1();
+  return;
+}
+

test/ptr_ref/store/param/ptr.spvasm.expected.spvasm

@@ -1,7 +1,7 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 19
+; Bound: 22
 ; Schema: 0
                OpCapability Shader
                OpMemoryModel Logical GLSL450
@@ -10,8 +10,9 @@
                OpName %func "func"
                OpName %value "value"
                OpName %pointer "pointer"
-               OpName %main "main"
+               OpName %main_1 "main_1"
                OpName %i "i"
+               OpName %main "main"
        %void = OpTypeVoid
         %int = OpTypeInt 32 1
 %_ptr_Function_int = OpTypePointer Function %int
@@ -27,7 +28,7 @@
                OpStore %pointer %value
                OpReturn
                OpFunctionEnd
-       %main = OpFunction %void None %10
+     %main_1 = OpFunction %void None %10
          %12 = OpLabel
           %i = OpVariable %_ptr_Function_int Function %15
                OpStore %i %int_0
@@ -35,3 +36,8 @@
          %17 = OpFunctionCall %void %func %int_123 %i
                OpReturn
                OpFunctionEnd
+       %main = OpFunction %void None %10
+         %20 = OpLabel
+         %21 = OpFunctionCall %void %main_1
+               OpReturn
+               OpFunctionEnd

test/ptr_ref/store/param/ptr.spvasm.expected.wgsl

@@ -3,10 +3,14 @@ fn func(value : i32, pointer : ptr<function, i32>) {
   return;
 }
 
-[[stage(compute)]]
-fn main() {
+fn main_1() {
   var i : i32 = 0;
   i = 123;
   func(123, &(i));
   return;
 }
+
+[[stage(compute)]]
+fn main() {
+  main_1();
+}

test/samples/simple_vertex.spvasm.expected.msl

@@ -1,13 +1,23 @@
 #include <metal_stdlib>
 
 using namespace metal;
-struct tint_symbol_out {
+struct main_out {
+  float4 gl_Position;
+};
+struct tint_symbol_1 {
   float4 gl_Position [[position]];
 };
 
-vertex tint_symbol_out tint_symbol() {
-  tint_symbol_out _tint_out = {};
-  _tint_out.gl_Position = float4(0.0f, 0.0f, 0.0f, 0.0f);
-  return _tint_out;
+void main_1(thread float4* const tint_symbol_4) {
+  *(tint_symbol_4) = float4(0.0f, 0.0f, 0.0f, 0.0f);
+  return;
+}
+
+vertex tint_symbol_1 tint_symbol() {
+  thread float4 tint_symbol_5 = 0.0f;
+  main_1(&(tint_symbol_5));
+  main_out const tint_symbol_2 = {.gl_Position=tint_symbol_5};
+  tint_symbol_1 const tint_symbol_3 = {.gl_Position=tint_symbol_2.gl_Position};
+  return tint_symbol_3;
 }
 

test/samples/simple_vertex.spvasm.expected.spvasm

@@ -1,32 +1,58 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 16
+; Bound: 30
 ; Schema: 0
                OpCapability Shader
                OpMemoryModel Logical GLSL450
-               OpEntryPoint Vertex %main "main" %tint_pointsize %gl_Position
+               OpEntryPoint Vertex %main "main" %tint_pointsize %tint_symbol_1
                OpName %tint_pointsize "tint_pointsize"
                OpName %gl_Position "gl_Position"
+               OpName %tint_symbol_1 "tint_symbol_1"
+               OpName %main_1 "main_1"
+               OpName %main_out "main_out"
+               OpMemberName %main_out 0 "gl_Position"
+               OpName %tint_symbol_2 "tint_symbol_2"
+               OpName %tint_symbol "tint_symbol"
                OpName %main "main"
                OpDecorate %tint_pointsize BuiltIn PointSize
-               OpDecorate %gl_Position BuiltIn Position
+               OpDecorate %tint_symbol_1 BuiltIn Position
+               OpMemberDecorate %main_out 0 Offset 0
       %float = OpTypeFloat 32
 %_ptr_Output_float = OpTypePointer Output %float
           %4 = OpConstantNull %float
 %tint_pointsize = OpVariable %_ptr_Output_float Output %4
     %v4float = OpTypeVector %float 4
-%_ptr_Output_v4float = OpTypePointer Output %v4float
+%_ptr_Private_v4float = OpTypePointer Private %v4float
           %8 = OpConstantNull %v4float
-%gl_Position = OpVariable %_ptr_Output_v4float Output %8
+%gl_Position = OpVariable %_ptr_Private_v4float Private %8
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+%tint_symbol_1 = OpVariable %_ptr_Output_v4float Output %8
        %void = OpTypeVoid
-          %9 = OpTypeFunction %void
-    %float_1 = OpConstant %float 1
+         %11 = OpTypeFunction %void
     %float_0 = OpConstant %float 0
-         %15 = OpConstantComposite %v4float %float_0 %float_0 %float_0 %float_0
-       %main = OpFunction %void None %9
-         %12 = OpLabel
-               OpStore %tint_pointsize %float_1
-               OpStore %gl_Position %15
+         %16 = OpConstantComposite %v4float %float_0 %float_0 %float_0 %float_0
+   %main_out = OpTypeStruct %v4float
+         %17 = OpTypeFunction %void %main_out
+    %float_1 = OpConstant %float 1
+     %main_1 = OpFunction %void None %11
+         %14 = OpLabel
+               OpStore %gl_Position %16
+               OpReturn
+               OpFunctionEnd
+%tint_symbol_2 = OpFunction %void None %17
+%tint_symbol = OpFunctionParameter %main_out
+         %21 = OpLabel
+         %22 = OpCompositeExtract %v4float %tint_symbol 0
+               OpStore %tint_symbol_1 %22
+               OpReturn
+               OpFunctionEnd
+       %main = OpFunction %void None %11
+         %24 = OpLabel
+               OpStore %tint_pointsize %float_1
+         %26 = OpFunctionCall %void %main_1
+         %28 = OpLoad %v4float %gl_Position
+         %29 = OpCompositeConstruct %main_out %28
+         %27 = OpFunctionCall %void %tint_symbol_2 %29
                OpReturn
                OpFunctionEnd

test/samples/simple_vertex.spvasm.expected.wgsl

@@ -1,7 +1,17 @@
-[[builtin(position)]] var<out> gl_Position : vec4<f32>;
+var<private> gl_Position : vec4<f32>;
 
-[[stage(vertex)]]
-fn main() {
+fn main_1() {
   gl_Position = vec4<f32>(0.0, 0.0, 0.0, 0.0);
   return;
 }
+
+struct main_out {
+  [[builtin(position)]]
+  gl_Position : vec4<f32>;
+};
+
+[[stage(vertex)]]
+fn main() -> main_out {
+  main_1();
+  return main_out(gl_Position);
+}