transform/msl: Run InlinePointerLets and Simplify

This will be relied on by the upcoming arrayLength transform.

Update test expectations.

Change-Id: Ib74b647abcd6f4393f9899ce40bbf06f6e53e7f4
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/55180
Commit-Queue: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>

src/transform/msl.cc

@@ -26,9 +26,11 @@
 #include "src/sem/variable.h"
 #include "src/transform/canonicalize_entry_point_io.h"
 #include "src/transform/external_texture_transform.h"
+#include "src/transform/inline_pointer_lets.h"
 #include "src/transform/manager.h"
 #include "src/transform/pad_array_elements.h"
 #include "src/transform/promote_initializers_to_const_var.h"
+#include "src/transform/simplify.h"
 #include "src/transform/wrap_arrays_in_structs.h"
 
 namespace tint {
@@ -45,6 +47,8 @@ Output Msl::Run(const Program* in, const DataMap&) {
   manager.Add<PromoteInitializersToConstVar>();
   manager.Add<WrapArraysInStructs>();
   manager.Add<PadArrayElements>();
+  manager.Add<InlinePointerLets>();
+  manager.Add<Simplify>();
   data.Add<CanonicalizeEntryPointIO::Config>(
       CanonicalizeEntryPointIO::BuiltinStyle::kParameter);
   auto out = manager.Run(in, data);

src/transform/msl_test.cc

@@ -144,10 +144,8 @@ fn main() {
 fn main() {
   [[internal(disable_validation__function_var_storage_class)]] var<private> tint_symbol : f32;
   [[internal(disable_validation__function_var_storage_class)]] var<workgroup> tint_symbol_1 : f32;
-  let p_ptr : ptr<private, f32> = &(tint_symbol);
+  let x : f32 = (tint_symbol + tint_symbol_1);
-  let w_ptr : ptr<workgroup, f32> = &(tint_symbol_1);
+  tint_symbol = x;
-  let x : f32 = (*(p_ptr) + *(w_ptr));
-  *(p_ptr) = x;
 }
 )";
 

test/array/assign_to_subexpr.wgsl.expected.msl

@@ -16,13 +16,10 @@ void foo() {
   tint_array_wrapper dst = {};
   S dst_struct = {};
   tint_array_wrapper_1 dst_array = {};
-  thread tint_array_wrapper* const dst_ptr = &(dst);
-  thread S* const dst_struct_ptr = &(dst_struct);
-  thread tint_array_wrapper_1* const dst_array_ptr = &(dst_array);
   dst_struct.arr = src;
   dst_array.arr[1] = src;
-  *(dst_ptr) = src;
+  dst = src;
-  (*(dst_struct_ptr)).arr = src;
+  dst_struct.arr = src;
-  (*(dst_array_ptr)).arr[0] = src;
+  dst_array.arr[0] = src;
 }
 

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

@@ -1,17 +1,19 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
+using namespace metal;
+void main_1(texture2d<uint, access::read> tint_symbol_1, texture2d<uint, access::write> tint_symbol_2) {
+  uint4 srcValue = 0u;
+  uint4 const x_18 = tint_symbol_1.read(uint2(int2(0, 0)));
+  srcValue = x_18;
+  uint const x_22 = srcValue.x;
+  srcValue.x = (x_22 + as_type<uint>(1));
+  uint4 const x_27 = srcValue;
+  tint_symbol_2.write(x_27, uint2(int2(0, 0)));
+  return;
+}
 
+kernel void tint_symbol(texture2d<uint, access::read> tint_symbol_3 [[texture(0)]], texture2d<uint, access::write> tint_symbol_4 [[texture(1)]]) {
+  main_1(tint_symbol_3, tint_symbol_4);
+  return;
+}
 
-Validation Failure:
-
-Compilation failed: 
-
-program_source:6:22: error: use of undeclared identifier 'Src'
-  uint4 const x_18 = Src.read(int2(0, 0));
-                     ^
-program_source:8:29: error: address of vector element requested
-  thread uint* const x_21 = &(srcValue.x);
-                            ^ ~~~~~~~~~~
-program_source:12:3: error: use of undeclared identifier 'Dst'
-  Dst.write(x_27, int2(0, 0));
-  ^

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

@@ -1,25 +1,14 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-bug/tint/453.wgsl:1:79 warning: use of deprecated language feature: access control is expected as last parameter of storage textures
+using namespace metal;
-[[group(0), binding(0)]] var Src : [[access(read)]] texture_storage_2d<r32uint>;
+kernel void tint_symbol(texture2d<uint, access::read> tint_symbol_1 [[texture(0)]], texture2d<uint, access::write> tint_symbol_2 [[texture(1)]]) {
-                                                                              ^
+  uint4 srcValue = 0u;
+  uint4 const x_22 = tint_symbol_1.read(uint2(int2(0, 0)));
+  srcValue = x_22;
+  uint const x_24 = srcValue.x;
+  uint const x_25 = (x_24 + 1u);
+  uint4 const x_27 = srcValue;
+  tint_symbol_2.write(x_27.xxxx, uint2(int2(0, 0)));
+  return;
+}
 
-bug/tint/453.wgsl:2:80 warning: use of deprecated language feature: access control is expected as last parameter of storage textures
-[[group(0), binding(1)]] var Dst : [[access(write)]] texture_storage_2d<r32uint>;
-                                                                               ^
-
-
-
-Validation Failure:
-
-Compilation failed: 
-
-program_source:6:22: error: use of undeclared identifier 'Src'
-  uint4 const x_22 = Src.read(int2(0, 0));
-                     ^
-program_source:8:29: error: address of vector element requested
-  thread uint* const x_23 = &(srcValue.x);
-                            ^ ~~~~~~~~~~
-program_source:12:3: error: use of undeclared identifier 'Dst'
-  Dst.write(x_27.xxxx, int2(0, 0));
-  ^

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

@@ -6,8 +6,7 @@ struct S {
 };
 
 kernel void tint_symbol(device S& buf [[buffer(0)]]) {
-  device int* const p = &(buf.a);
+  buf.a = 12;
-  *(p) = 12;
   return;
 }
 

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

@@ -3,8 +3,7 @@
 using namespace metal;
 kernel void tint_symbol() {
   float3x3 m = float3x3(float3(1.0f, 2.0f, 3.0f), float3(4.0f, 5.0f, 6.0f), float3(7.0f, 8.0f, 9.0f));
-  thread float3* const v = &(m[1]);
+  m[1] = float3(5.0f, 5.0f, 5.0f);
-  *(v) = float3(5.0f, 5.0f, 5.0f);
   return;
 }
 

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

@@ -1,12 +1,9 @@
-SKIP: crbug.com/tint/816
-
 #include <metal_stdlib>
 
 using namespace metal;
 kernel void tint_symbol() {
   float3 v = float3(1.0f, 2.0f, 3.0f);
-  thread float* const f = &(v.y);
+  v.y = 5.0f;
-  *(f) = 5.0f;
   return;
 }
 

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

@@ -3,8 +3,6 @@
 using namespace metal;
 void main_1() {
   uint x_10 = 0u;
-  thread uint* const x_1 = &(x_10);
-  thread uint* const x_2 = x_1;
   return;
 }
 

test/ptr_ref/load/local/ptr_function.wgsl.expected.msl

@@ -3,8 +3,7 @@
 using namespace metal;
 kernel void tint_symbol() {
   int i = 123;
-  thread int* const p = &(i);
+  int const use = (i + 1);
-  int const use = (*(p) + 1);
   return;
 }
 

test/ptr_ref/load/local/ptr_private.wgsl.expected.msl

@@ -3,8 +3,7 @@
 using namespace metal;
 kernel void tint_symbol() {
   thread int tint_symbol_1 = 123;
-  thread int* const p = &(tint_symbol_1);
+  int const use = (tint_symbol_1 + 1);
-  int const use = (*(p) + 1);
   return;
 }
 

test/ptr_ref/load/local/ptr_storage.wgsl.expected.msl

@@ -6,8 +6,7 @@ struct S {
 };
 
 kernel void tint_symbol(device S& v [[buffer(0)]]) {
-  device int* const p = &(v.a);
+  int const use = (v.a + 1);
-  int const use = (*(p) + 1);
   return;
 }
 

test/ptr_ref/load/local/ptr_uniform.wgsl.expected.msl

@@ -6,8 +6,7 @@ struct S {
 };
 
 kernel void tint_symbol(constant S& v [[buffer(0)]]) {
-  constant int* const p = &(v.a);
+  int const use = (v.a + 1);
-  int const use = (*(p) + 1);
   return;
 }
 

test/ptr_ref/load/local/ptr_workgroup.wgsl.expected.msl

@@ -4,8 +4,7 @@ using namespace metal;
 kernel void tint_symbol() {
   threadgroup int tint_symbol_1 = 0;
   tint_symbol_1 = 123;
-  threadgroup int* const p = &(tint_symbol_1);
+  int const use = (tint_symbol_1 + 1);
-  int const use = (*(p) + 1);
   return;
 }
 

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

@@ -3,9 +3,8 @@
 using namespace metal;
 kernel void tint_symbol() {
   int i = 123;
-  thread int* const p = &(i);
+  i = 123;
-  *(p) = 123;
+  i = ((100 + 20) + 3);
-  *(p) = ((100 + 20) + 3);
   return;
 }
 

test/types/function_scope_declarations.wgsl.expected.msl

@@ -28,9 +28,6 @@ kernel void tint_symbol() {
   tint_array_wrapper const arr_let = {.arr={}};
   S struct_var = {};
   S const struct_let = {};
-  thread float* const ptr_f32 = &(f32_var);
-  thread float4* const ptr_vec = &(v4f32_var);
-  thread tint_array_wrapper* const ptr_arr = &(arr_var);
   return;
 }