From 567f2e4f3b084c602db781f263e4062d64aaaffa Mon Sep 17 00:00:00 2001 From: James Price Date: Fri, 18 Jun 2021 09:47:23 +0000 Subject: [PATCH] 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 Kokoro: Kokoro Reviewed-by: Ben Clayton --- src/transform/msl.cc | 4 +++ src/transform/msl_test.cc | 6 ++-- .../array/assign_to_subexpr.wgsl.expected.msl | 9 ++--- test/bug/tint/413.spvasm.expected.msl | 30 ++++++++-------- test/bug/tint/453.wgsl.expected.msl | 35 +++++++------------ test/bug/tint/492.wgsl.expected.msl | 3 +- test/ptr_ref/access/matrix.wgsl.expected.msl | 3 +- test/ptr_ref/access/vector.wgsl.expected.msl | 5 +-- .../ptr_ref/copy/ptr_copy.spvasm.expected.msl | 2 -- .../load/local/ptr_function.wgsl.expected.msl | 3 +- .../load/local/ptr_private.wgsl.expected.msl | 3 +- .../load/local/ptr_storage.wgsl.expected.msl | 3 +- .../load/local/ptr_uniform.wgsl.expected.msl | 3 +- .../local/ptr_workgroup.wgsl.expected.msl | 3 +- .../ptr_ref/store/local/i32.wgsl.expected.msl | 5 ++- ...ction_scope_declarations.wgsl.expected.msl | 3 -- 16 files changed, 47 insertions(+), 73 deletions(-) diff --git a/src/transform/msl.cc b/src/transform/msl.cc index 8b73c422b8..10e3ea27ec 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 @@ Output Msl::Run(const Program* in, const DataMap&) { manager.Add(); manager.Add(); manager.Add(); + manager.Add(); + manager.Add(); data.Add( CanonicalizeEntryPointIO::BuiltinStyle::kParameter); auto out = manager.Run(in, data); diff --git a/src/transform/msl_test.cc b/src/transform/msl_test.cc index 21553f2821..6e19c44ef7 100644 --- a/src/transform/msl_test.cc +++ b/src/transform/msl_test.cc @@ -144,10 +144,8 @@ fn main() { fn main() { [[internal(disable_validation__function_var_storage_class)]] var tint_symbol : f32; [[internal(disable_validation__function_var_storage_class)]] var tint_symbol_1 : f32; - let p_ptr : ptr = &(tint_symbol); - let w_ptr : ptr = &(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 2b62433a6b..9b4779214c 100644 --- a/test/array/assign_to_subexpr.wgsl.expected.msl +++ b/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_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 ad93adb873..ac3858a3ef 100644 --- a/test/bug/tint/413.spvasm.expected.msl +++ b/test/bug/tint/413.spvasm.expected.msl @@ -1,17 +1,19 @@ -SKIP: FAILED +#include +using namespace metal; +void main_1(texture2d tint_symbol_1, texture2d 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(1)); + uint4 const x_27 = srcValue; + tint_symbol_2.write(x_27, uint2(int2(0, 0))); + return; +} +kernel void tint_symbol(texture2d tint_symbol_3 [[texture(0)]], texture2d 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 fc9acbed1f..e5de3d4433 100644 --- a/test/bug/tint/453.wgsl.expected.msl +++ b/test/bug/tint/453.wgsl.expected.msl @@ -1,25 +1,14 @@ -SKIP: FAILED +#include -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; - ^ +using namespace metal; +kernel void tint_symbol(texture2d tint_symbol_1 [[texture(0)]], texture2d 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; - ^ - - - -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 ddf883fb03..b7caff157d 100644 --- a/test/bug/tint/492.wgsl.expected.msl +++ b/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); - *(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 1a95911ba5..75c13af6bb 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 87122cf971..6625cdb8d1 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 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 5495c634ee..ff7940b89d 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 ce60e0c8c5..e7b87bb817 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 c600e5b915..5ceee0ad2a 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 17d0c1bf6f..c00663c41e 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 @@ struct S { }; 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 9008b2fbd3..5aeb08155e 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 @@ struct S { }; 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 1f4dfe844e..4c51f84b40 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 @@ 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 = (*(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 3c243ed285..bde28674c1 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 3587c458c9..0c50b0227d 100644 --- a/test/types/function_scope_declarations.wgsl.expected.msl +++ b/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; }