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>
diff --git a/src/transform/msl.cc b/src/transform/msl.cc
index 8b73c42..10e3ea2 100644
--- a/src/transform/msl.cc
+++ b/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 @@
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);
diff --git a/src/transform/msl_test.cc b/src/transform/msl_test.cc
index 21553f2..6e19c44 100644
--- a/src/transform/msl_test.cc
+++ b/src/transform/msl_test.cc
@@ -144,10 +144,8 @@
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 w_ptr : ptr<workgroup, f32> = &(tint_symbol_1);
- let x : f32 = (*(p_ptr) + *(w_ptr));
- *(p_ptr) = x;
+ let x : f32 = (tint_symbol + tint_symbol_1);
+ tint_symbol = x;
}
)";
diff --git a/test/array/assign_to_subexpr.wgsl.expected.msl b/test/array/assign_to_subexpr.wgsl.expected.msl
index 2b62433..9b47792 100644
--- a/test/array/assign_to_subexpr.wgsl.expected.msl
+++ b/test/array/assign_to_subexpr.wgsl.expected.msl
@@ -16,13 +16,10 @@
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_struct_ptr)).arr = src;
- (*(dst_array_ptr)).arr[0] = src;
+ dst = src;
+ dst_struct.arr = src;
+ dst_array.arr[0] = src;
}
diff --git a/test/bug/tint/413.spvasm.expected.msl b/test/bug/tint/413.spvasm.expected.msl
index ad93adb..ac3858a 100644
--- a/test/bug/tint/413.spvasm.expected.msl
+++ b/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));
- ^
diff --git a/test/bug/tint/453.wgsl.expected.msl b/test/bug/tint/453.wgsl.expected.msl
index fc9acbe..e5de3d4 100644
--- a/test/bug/tint/453.wgsl.expected.msl
+++ b/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
-[[group(0), binding(0)]] var Src : [[access(read)]] texture_storage_2d<r32uint>;
- ^
+using namespace metal;
+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));
- ^
diff --git a/test/bug/tint/492.wgsl.expected.msl b/test/bug/tint/492.wgsl.expected.msl
index ddf883f..b7caff1 100644
--- a/test/bug/tint/492.wgsl.expected.msl
+++ b/test/bug/tint/492.wgsl.expected.msl
@@ -6,8 +6,7 @@
};
kernel void tint_symbol(device S& buf [[buffer(0)]]) {
- device int* const p = &(buf.a);
- *(p) = 12;
+ buf.a = 12;
return;
}
diff --git a/test/ptr_ref/access/matrix.wgsl.expected.msl b/test/ptr_ref/access/matrix.wgsl.expected.msl
index 1a95911..75c13af 100644
--- a/test/ptr_ref/access/matrix.wgsl.expected.msl
+++ b/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]);
- *(v) = float3(5.0f, 5.0f, 5.0f);
+ m[1] = float3(5.0f, 5.0f, 5.0f);
return;
}
diff --git a/test/ptr_ref/access/vector.wgsl.expected.msl b/test/ptr_ref/access/vector.wgsl.expected.msl
index 87122cf..6625cdb 100644
--- a/test/ptr_ref/access/vector.wgsl.expected.msl
+++ b/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);
- *(f) = 5.0f;
+ v.y = 5.0f;
return;
}
diff --git a/test/ptr_ref/copy/ptr_copy.spvasm.expected.msl b/test/ptr_ref/copy/ptr_copy.spvasm.expected.msl
index 5495c63..ff7940b 100644
--- a/test/ptr_ref/copy/ptr_copy.spvasm.expected.msl
+++ b/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;
}
diff --git a/test/ptr_ref/load/local/ptr_function.wgsl.expected.msl b/test/ptr_ref/load/local/ptr_function.wgsl.expected.msl
index ce60e0c..e7b87bb 100644
--- a/test/ptr_ref/load/local/ptr_function.wgsl.expected.msl
+++ b/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 = (*(p) + 1);
+ int const use = (i + 1);
return;
}
diff --git a/test/ptr_ref/load/local/ptr_private.wgsl.expected.msl b/test/ptr_ref/load/local/ptr_private.wgsl.expected.msl
index c600e5b..5ceee0a 100644
--- a/test/ptr_ref/load/local/ptr_private.wgsl.expected.msl
+++ b/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 = (*(p) + 1);
+ int const use = (tint_symbol_1 + 1);
return;
}
diff --git a/test/ptr_ref/load/local/ptr_storage.wgsl.expected.msl b/test/ptr_ref/load/local/ptr_storage.wgsl.expected.msl
index 17d0c1b..c00663c 100644
--- a/test/ptr_ref/load/local/ptr_storage.wgsl.expected.msl
+++ b/test/ptr_ref/load/local/ptr_storage.wgsl.expected.msl
@@ -6,8 +6,7 @@
};
kernel void tint_symbol(device S& v [[buffer(0)]]) {
- device int* const p = &(v.a);
- int const use = (*(p) + 1);
+ int const use = (v.a + 1);
return;
}
diff --git a/test/ptr_ref/load/local/ptr_uniform.wgsl.expected.msl b/test/ptr_ref/load/local/ptr_uniform.wgsl.expected.msl
index 9008b2f..5aeb081 100644
--- a/test/ptr_ref/load/local/ptr_uniform.wgsl.expected.msl
+++ b/test/ptr_ref/load/local/ptr_uniform.wgsl.expected.msl
@@ -6,8 +6,7 @@
};
kernel void tint_symbol(constant S& v [[buffer(0)]]) {
- constant int* const p = &(v.a);
- int const use = (*(p) + 1);
+ int const use = (v.a + 1);
return;
}
diff --git a/test/ptr_ref/load/local/ptr_workgroup.wgsl.expected.msl b/test/ptr_ref/load/local/ptr_workgroup.wgsl.expected.msl
index 1f4dfe8..4c51f84 100644
--- a/test/ptr_ref/load/local/ptr_workgroup.wgsl.expected.msl
+++ b/test/ptr_ref/load/local/ptr_workgroup.wgsl.expected.msl
@@ -4,8 +4,7 @@
kernel void tint_symbol() {
threadgroup int tint_symbol_1 = 0;
tint_symbol_1 = 123;
- threadgroup int* const p = &(tint_symbol_1);
- int const use = (*(p) + 1);
+ int const use = (tint_symbol_1 + 1);
return;
}
diff --git a/test/ptr_ref/store/local/i32.wgsl.expected.msl b/test/ptr_ref/store/local/i32.wgsl.expected.msl
index 3c243ed..bde2867 100644
--- a/test/ptr_ref/store/local/i32.wgsl.expected.msl
+++ b/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);
- *(p) = 123;
- *(p) = ((100 + 20) + 3);
+ i = 123;
+ i = ((100 + 20) + 3);
return;
}
diff --git a/test/types/function_scope_declarations.wgsl.expected.msl b/test/types/function_scope_declarations.wgsl.expected.msl
index 3587c45..0c50b02 100644
--- a/test/types/function_scope_declarations.wgsl.expected.msl
+++ b/test/types/function_scope_declarations.wgsl.expected.msl
@@ -28,9 +28,6 @@
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;
}
