tint: add syntax sugar for dereferencing pointers
Adds support to Tint for syntax sugar for dereferencing pointers for
member or index access as per:
https://github.com/gpuweb/gpuweb/pull/4311
- Resolver: when the lhs of a accessor expression is a pointer, it is
now resolved to a sem::Reference.
- Added "pointer_composite_access" feature as experimental, hooked up
validation in Resolver, and added tests.
- Added resolver tests for the new syntax to resolver/ptr_ref_test.cc.
- Fixed multiple transforms to deal with the fact that the lhs of
accessor expressions can now be pointers, including: Robustness,
Renamer, Std140, and SimplifyPointers.
- In transforms that rely on other transforms, such as
SimplifyPointers, to remove/inline pointers, I added asserts that
the type is not a pointer.
- Added unit tests for transforms that use pointer-dot/index for
accessor expressions.
- Fixed uniformity analysis code so that ProcessLValueExpression
correctly deals with accessor expressions where the object is a
pointer, in the same way we do for UnaryOp::kIndirection, including
partial pointer checks. Added many tests for these new cases.
- Fixed ProgramToIR so that EmitAccess handles the new syntax. Added
multiple tests.
- Added end2end tests under test/tint/ptr_sugar
For Googlers, see my work log at
go/add-syntax-sugar-for-dereferencing-composites for more details.
Bug: tint:2113
Change-Id: I7a0093f52ca2237be598e44245b45049f21d056c
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/164900
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Corentin Wallez <cwallez@chromium.org>
Reviewed-by: James Price <jrprice@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
Commit-Queue: Antonio Maiorano <amaiorano@google.com>
diff --git a/test/tint/ptr_sugar/array.wgsl b/test/tint/ptr_sugar/array.wgsl
new file mode 100644
index 0000000..d48a62b
--- /dev/null
+++ b/test/tint/ptr_sugar/array.wgsl
@@ -0,0 +1,55 @@
+fn deref_const() {
+ var a : array<i32, 10>;
+ let p = &a;
+ var b = (*p)[0];
+ (*p)[0] = 42;
+}
+
+fn no_deref_const() {
+ var a : array<i32, 10>;
+ let p = &a;
+ var b = p[0];
+ p[0] = 42;
+}
+
+fn deref_let() {
+ var a : array<i32, 10>;
+ let p = &a;
+ let i = 0;
+ var b = (*p)[i];
+ (*p)[0] = 42;
+}
+
+fn no_deref_let() {
+ var a : array<i32, 10>;
+ let p = &a;
+ let i = 0;
+ var b = p[i];
+ p[0] = 42;
+}
+
+fn deref_var() {
+ var a : array<i32, 10>;
+ let p = &a;
+ var i = 0;
+ var b = (*p)[i];
+ (*p)[0] = 42;
+}
+
+fn no_deref_var() {
+ var a : array<i32, 10>;
+ let p = &a;
+ var i = 0;
+ var b = p[i];
+ p[0] = 42;
+}
+
+@compute @workgroup_size(1)
+fn main() {
+ deref_const();
+ no_deref_const();
+ deref_let();
+ no_deref_let();
+ deref_var();
+ no_deref_var();
+}
diff --git a/test/tint/ptr_sugar/array.wgsl.expected.dxc.hlsl b/test/tint/ptr_sugar/array.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..c876ea4
--- /dev/null
+++ b/test/tint/ptr_sugar/array.wgsl.expected.dxc.hlsl
@@ -0,0 +1,50 @@
+void deref_const() {
+ int a[10] = (int[10])0;
+ int b = a[0];
+ a[0] = 42;
+}
+
+void no_deref_const() {
+ int a[10] = (int[10])0;
+ int b = a[0];
+ a[0] = 42;
+}
+
+void deref_let() {
+ int a[10] = (int[10])0;
+ const int i = 0;
+ int b = a[i];
+ a[0] = 42;
+}
+
+void no_deref_let() {
+ int a[10] = (int[10])0;
+ const int i = 0;
+ int b = a[i];
+ a[0] = 42;
+}
+
+void deref_var() {
+ int a[10] = (int[10])0;
+ int i = 0;
+ int b = a[i];
+ a[0] = 42;
+}
+
+void no_deref_var() {
+ int a[10] = (int[10])0;
+ int i = 0;
+ int b = a[i];
+ a[0] = 42;
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+ deref_const();
+ no_deref_const();
+ deref_let();
+ no_deref_let();
+ deref_var();
+ no_deref_var();
+ return;
+}
diff --git a/test/tint/ptr_sugar/array.wgsl.expected.fxc.hlsl b/test/tint/ptr_sugar/array.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..c876ea4
--- /dev/null
+++ b/test/tint/ptr_sugar/array.wgsl.expected.fxc.hlsl
@@ -0,0 +1,50 @@
+void deref_const() {
+ int a[10] = (int[10])0;
+ int b = a[0];
+ a[0] = 42;
+}
+
+void no_deref_const() {
+ int a[10] = (int[10])0;
+ int b = a[0];
+ a[0] = 42;
+}
+
+void deref_let() {
+ int a[10] = (int[10])0;
+ const int i = 0;
+ int b = a[i];
+ a[0] = 42;
+}
+
+void no_deref_let() {
+ int a[10] = (int[10])0;
+ const int i = 0;
+ int b = a[i];
+ a[0] = 42;
+}
+
+void deref_var() {
+ int a[10] = (int[10])0;
+ int i = 0;
+ int b = a[i];
+ a[0] = 42;
+}
+
+void no_deref_var() {
+ int a[10] = (int[10])0;
+ int i = 0;
+ int b = a[i];
+ a[0] = 42;
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+ deref_const();
+ no_deref_const();
+ deref_let();
+ no_deref_let();
+ deref_var();
+ no_deref_var();
+ return;
+}
diff --git a/test/tint/ptr_sugar/array.wgsl.expected.glsl b/test/tint/ptr_sugar/array.wgsl.expected.glsl
new file mode 100644
index 0000000..48419fc
--- /dev/null
+++ b/test/tint/ptr_sugar/array.wgsl.expected.glsl
@@ -0,0 +1,56 @@
+#version 310 es
+
+void deref_const() {
+ int a[10] = int[10](0, 0, 0, 0, 0, 0, 0, 0, 0, 0);
+ int b = a[0];
+ a[0] = 42;
+}
+
+void no_deref_const() {
+ int a[10] = int[10](0, 0, 0, 0, 0, 0, 0, 0, 0, 0);
+ int b = a[0];
+ a[0] = 42;
+}
+
+void deref_let() {
+ int a[10] = int[10](0, 0, 0, 0, 0, 0, 0, 0, 0, 0);
+ int i = 0;
+ int b = a[i];
+ a[0] = 42;
+}
+
+void no_deref_let() {
+ int a[10] = int[10](0, 0, 0, 0, 0, 0, 0, 0, 0, 0);
+ int i = 0;
+ int b = a[i];
+ a[0] = 42;
+}
+
+void deref_var() {
+ int a[10] = int[10](0, 0, 0, 0, 0, 0, 0, 0, 0, 0);
+ int i = 0;
+ int b = a[i];
+ a[0] = 42;
+}
+
+void no_deref_var() {
+ int a[10] = int[10](0, 0, 0, 0, 0, 0, 0, 0, 0, 0);
+ int i = 0;
+ int b = a[i];
+ a[0] = 42;
+}
+
+void tint_symbol() {
+ deref_const();
+ no_deref_const();
+ deref_let();
+ no_deref_let();
+ deref_var();
+ no_deref_var();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ tint_symbol();
+ return;
+}
diff --git a/test/tint/ptr_sugar/array.wgsl.expected.msl b/test/tint/ptr_sugar/array.wgsl.expected.msl
new file mode 100644
index 0000000..9bd0861
--- /dev/null
+++ b/test/tint/ptr_sugar/array.wgsl.expected.msl
@@ -0,0 +1,66 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
+};
+
+void deref_const() {
+ tint_array<int, 10> a = {};
+ int b = a[0];
+ a[0] = 42;
+}
+
+void no_deref_const() {
+ tint_array<int, 10> a = {};
+ int b = a[0];
+ a[0] = 42;
+}
+
+void deref_let() {
+ tint_array<int, 10> a = {};
+ int const i = 0;
+ int b = a[i];
+ a[0] = 42;
+}
+
+void no_deref_let() {
+ tint_array<int, 10> a = {};
+ int const i = 0;
+ int b = a[i];
+ a[0] = 42;
+}
+
+void deref_var() {
+ tint_array<int, 10> a = {};
+ int i = 0;
+ int b = a[i];
+ a[0] = 42;
+}
+
+void no_deref_var() {
+ tint_array<int, 10> a = {};
+ int i = 0;
+ int b = a[i];
+ a[0] = 42;
+}
+
+kernel void tint_symbol() {
+ deref_const();
+ no_deref_const();
+ deref_let();
+ no_deref_let();
+ deref_var();
+ no_deref_var();
+ return;
+}
+
diff --git a/test/tint/ptr_sugar/array.wgsl.expected.spvasm b/test/tint/ptr_sugar/array.wgsl.expected.spvasm
new file mode 100644
index 0000000..fd7b149
--- /dev/null
+++ b/test/tint/ptr_sugar/array.wgsl.expected.spvasm
@@ -0,0 +1,124 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 66
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %deref_const "deref_const"
+ OpName %a "a"
+ OpName %b "b"
+ OpName %no_deref_const "no_deref_const"
+ OpName %a_0 "a"
+ OpName %b_0 "b"
+ OpName %deref_let "deref_let"
+ OpName %a_1 "a"
+ OpName %b_1 "b"
+ OpName %no_deref_let "no_deref_let"
+ OpName %a_2 "a"
+ OpName %b_2 "b"
+ OpName %deref_var "deref_var"
+ OpName %a_3 "a"
+ OpName %i "i"
+ OpName %b_3 "b"
+ OpName %no_deref_var "no_deref_var"
+ OpName %a_4 "a"
+ OpName %i_0 "i"
+ OpName %b_4 "b"
+ OpName %main "main"
+ OpDecorate %_arr_int_uint_10 ArrayStride 4
+ %void = OpTypeVoid
+ %1 = OpTypeFunction %void
+ %int = OpTypeInt 32 1
+ %uint = OpTypeInt 32 0
+ %uint_10 = OpConstant %uint 10
+%_arr_int_uint_10 = OpTypeArray %int %uint_10
+%_ptr_Function__arr_int_uint_10 = OpTypePointer Function %_arr_int_uint_10
+ %11 = OpConstantNull %_arr_int_uint_10
+ %12 = OpConstantNull %int
+%_ptr_Function_int = OpTypePointer Function %int
+ %int_42 = OpConstant %int 42
+%deref_const = OpFunction %void None %1
+ %4 = OpLabel
+ %a = OpVariable %_ptr_Function__arr_int_uint_10 Function %11
+ %b = OpVariable %_ptr_Function_int Function %12
+ %14 = OpAccessChain %_ptr_Function_int %a %12
+ %15 = OpLoad %int %14
+ OpStore %b %15
+ %17 = OpAccessChain %_ptr_Function_int %a %12
+ OpStore %17 %int_42
+ OpReturn
+ OpFunctionEnd
+%no_deref_const = OpFunction %void None %1
+ %20 = OpLabel
+ %a_0 = OpVariable %_ptr_Function__arr_int_uint_10 Function %11
+ %b_0 = OpVariable %_ptr_Function_int Function %12
+ %22 = OpAccessChain %_ptr_Function_int %a_0 %12
+ %23 = OpLoad %int %22
+ OpStore %b_0 %23
+ %25 = OpAccessChain %_ptr_Function_int %a_0 %12
+ OpStore %25 %int_42
+ OpReturn
+ OpFunctionEnd
+ %deref_let = OpFunction %void None %1
+ %27 = OpLabel
+ %a_1 = OpVariable %_ptr_Function__arr_int_uint_10 Function %11
+ %b_1 = OpVariable %_ptr_Function_int Function %12
+ %29 = OpAccessChain %_ptr_Function_int %a_1 %12
+ %30 = OpLoad %int %29
+ OpStore %b_1 %30
+ %32 = OpAccessChain %_ptr_Function_int %a_1 %12
+ OpStore %32 %int_42
+ OpReturn
+ OpFunctionEnd
+%no_deref_let = OpFunction %void None %1
+ %34 = OpLabel
+ %a_2 = OpVariable %_ptr_Function__arr_int_uint_10 Function %11
+ %b_2 = OpVariable %_ptr_Function_int Function %12
+ %36 = OpAccessChain %_ptr_Function_int %a_2 %12
+ %37 = OpLoad %int %36
+ OpStore %b_2 %37
+ %39 = OpAccessChain %_ptr_Function_int %a_2 %12
+ OpStore %39 %int_42
+ OpReturn
+ OpFunctionEnd
+ %deref_var = OpFunction %void None %1
+ %41 = OpLabel
+ %a_3 = OpVariable %_ptr_Function__arr_int_uint_10 Function %11
+ %i = OpVariable %_ptr_Function_int Function %12
+ %b_3 = OpVariable %_ptr_Function_int Function %12
+ OpStore %i %12
+ %44 = OpLoad %int %i
+ %45 = OpAccessChain %_ptr_Function_int %a_3 %44
+ %46 = OpLoad %int %45
+ OpStore %b_3 %46
+ %48 = OpAccessChain %_ptr_Function_int %a_3 %12
+ OpStore %48 %int_42
+ OpReturn
+ OpFunctionEnd
+%no_deref_var = OpFunction %void None %1
+ %50 = OpLabel
+ %a_4 = OpVariable %_ptr_Function__arr_int_uint_10 Function %11
+ %i_0 = OpVariable %_ptr_Function_int Function %12
+ %b_4 = OpVariable %_ptr_Function_int Function %12
+ OpStore %i_0 %12
+ %53 = OpLoad %int %i_0
+ %54 = OpAccessChain %_ptr_Function_int %a_4 %53
+ %55 = OpLoad %int %54
+ OpStore %b_4 %55
+ %57 = OpAccessChain %_ptr_Function_int %a_4 %12
+ OpStore %57 %int_42
+ OpReturn
+ OpFunctionEnd
+ %main = OpFunction %void None %1
+ %59 = OpLabel
+ %60 = OpFunctionCall %void %deref_const
+ %61 = OpFunctionCall %void %no_deref_const
+ %62 = OpFunctionCall %void %deref_let
+ %63 = OpFunctionCall %void %no_deref_let
+ %64 = OpFunctionCall %void %deref_var
+ %65 = OpFunctionCall %void %no_deref_var
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/ptr_sugar/array.wgsl.expected.wgsl b/test/tint/ptr_sugar/array.wgsl.expected.wgsl
new file mode 100644
index 0000000..da29411
--- /dev/null
+++ b/test/tint/ptr_sugar/array.wgsl.expected.wgsl
@@ -0,0 +1,55 @@
+fn deref_const() {
+ var a : array<i32, 10>;
+ let p = &(a);
+ var b = (*(p))[0];
+ (*(p))[0] = 42;
+}
+
+fn no_deref_const() {
+ var a : array<i32, 10>;
+ let p = &(a);
+ var b = p[0];
+ p[0] = 42;
+}
+
+fn deref_let() {
+ var a : array<i32, 10>;
+ let p = &(a);
+ let i = 0;
+ var b = (*(p))[i];
+ (*(p))[0] = 42;
+}
+
+fn no_deref_let() {
+ var a : array<i32, 10>;
+ let p = &(a);
+ let i = 0;
+ var b = p[i];
+ p[0] = 42;
+}
+
+fn deref_var() {
+ var a : array<i32, 10>;
+ let p = &(a);
+ var i = 0;
+ var b = (*(p))[i];
+ (*(p))[0] = 42;
+}
+
+fn no_deref_var() {
+ var a : array<i32, 10>;
+ let p = &(a);
+ var i = 0;
+ var b = p[i];
+ p[0] = 42;
+}
+
+@compute @workgroup_size(1)
+fn main() {
+ deref_const();
+ no_deref_const();
+ deref_let();
+ no_deref_let();
+ deref_var();
+ no_deref_var();
+}
diff --git a/test/tint/ptr_sugar/builtin_struct.wgsl b/test/tint/ptr_sugar/builtin_struct.wgsl
new file mode 100644
index 0000000..b44cbc4
--- /dev/null
+++ b/test/tint/ptr_sugar/builtin_struct.wgsl
@@ -0,0 +1,35 @@
+fn deref_modf() {
+ var a = modf(1.5);
+ let p = &a;
+ var fract = (*p).fract;
+ var whole = (*p).whole;
+}
+
+fn no_deref_modf() {
+ var a = modf(1.5);
+ let p = &a;
+ var fract = p.fract;
+ var whole = p.whole;
+}
+
+fn deref_frexp() {
+ var a = frexp(1.5);
+ let p = &a;
+ var fract = (*p).fract;
+ var exp = (*p).exp;
+}
+
+fn no_deref_frexp() {
+ var a = frexp(1.5);
+ let p = &a;
+ var fract = p.fract;
+ var exp = p.exp;
+}
+
+@compute @workgroup_size(1)
+fn main() {
+ deref_modf();
+ no_deref_modf();
+ deref_frexp();
+ no_deref_frexp();
+}
diff --git a/test/tint/ptr_sugar/builtin_struct.wgsl.expected.dxc.hlsl b/test/tint/ptr_sugar/builtin_struct.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..ec96b03
--- /dev/null
+++ b/test/tint/ptr_sugar/builtin_struct.wgsl.expected.dxc.hlsl
@@ -0,0 +1,40 @@
+struct modf_result_f32 {
+ float fract;
+ float whole;
+};
+struct frexp_result_f32 {
+ float fract;
+ int exp;
+};
+void deref_modf() {
+ modf_result_f32 a = {0.5f, 1.0f};
+ float fract = a.fract;
+ float whole = a.whole;
+}
+
+void no_deref_modf() {
+ modf_result_f32 a = {0.5f, 1.0f};
+ float fract = a.fract;
+ float whole = a.whole;
+}
+
+void deref_frexp() {
+ frexp_result_f32 a = {0.75f, 1};
+ float fract = a.fract;
+ int exp = a.exp;
+}
+
+void no_deref_frexp() {
+ frexp_result_f32 a = {0.75f, 1};
+ float fract = a.fract;
+ int exp = a.exp;
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+ deref_modf();
+ no_deref_modf();
+ deref_frexp();
+ no_deref_frexp();
+ return;
+}
diff --git a/test/tint/ptr_sugar/builtin_struct.wgsl.expected.fxc.hlsl b/test/tint/ptr_sugar/builtin_struct.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..ec96b03
--- /dev/null
+++ b/test/tint/ptr_sugar/builtin_struct.wgsl.expected.fxc.hlsl
@@ -0,0 +1,40 @@
+struct modf_result_f32 {
+ float fract;
+ float whole;
+};
+struct frexp_result_f32 {
+ float fract;
+ int exp;
+};
+void deref_modf() {
+ modf_result_f32 a = {0.5f, 1.0f};
+ float fract = a.fract;
+ float whole = a.whole;
+}
+
+void no_deref_modf() {
+ modf_result_f32 a = {0.5f, 1.0f};
+ float fract = a.fract;
+ float whole = a.whole;
+}
+
+void deref_frexp() {
+ frexp_result_f32 a = {0.75f, 1};
+ float fract = a.fract;
+ int exp = a.exp;
+}
+
+void no_deref_frexp() {
+ frexp_result_f32 a = {0.75f, 1};
+ float fract = a.fract;
+ int exp = a.exp;
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+ deref_modf();
+ no_deref_modf();
+ deref_frexp();
+ no_deref_frexp();
+ return;
+}
diff --git a/test/tint/ptr_sugar/builtin_struct.wgsl.expected.glsl b/test/tint/ptr_sugar/builtin_struct.wgsl.expected.glsl
new file mode 100644
index 0000000..e852152
--- /dev/null
+++ b/test/tint/ptr_sugar/builtin_struct.wgsl.expected.glsl
@@ -0,0 +1,49 @@
+#version 310 es
+
+struct modf_result_f32 {
+ float fract;
+ float whole;
+};
+
+struct frexp_result_f32 {
+ float fract;
+ int exp;
+};
+
+
+void deref_modf() {
+ modf_result_f32 a = modf_result_f32(0.5f, 1.0f);
+ float tint_symbol = a.fract;
+ float whole = a.whole;
+}
+
+void no_deref_modf() {
+ modf_result_f32 a = modf_result_f32(0.5f, 1.0f);
+ float tint_symbol = a.fract;
+ float whole = a.whole;
+}
+
+void deref_frexp() {
+ frexp_result_f32 a = frexp_result_f32(0.75f, 1);
+ float tint_symbol = a.fract;
+ int tint_symbol_1 = a.exp;
+}
+
+void no_deref_frexp() {
+ frexp_result_f32 a = frexp_result_f32(0.75f, 1);
+ float tint_symbol = a.fract;
+ int tint_symbol_1 = a.exp;
+}
+
+void tint_symbol_2() {
+ deref_modf();
+ no_deref_modf();
+ deref_frexp();
+ no_deref_frexp();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ tint_symbol_2();
+ return;
+}
diff --git a/test/tint/ptr_sugar/builtin_struct.wgsl.expected.msl b/test/tint/ptr_sugar/builtin_struct.wgsl.expected.msl
new file mode 100644
index 0000000..13aa15f
--- /dev/null
+++ b/test/tint/ptr_sugar/builtin_struct.wgsl.expected.msl
@@ -0,0 +1,44 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+struct modf_result_f32 {
+ float fract;
+ float whole;
+};
+struct frexp_result_f32 {
+ float fract;
+ int exp;
+};
+void deref_modf() {
+ modf_result_f32 a = modf_result_f32{.fract=0.5f, .whole=1.0f};
+ float fract = a.fract;
+ float whole = a.whole;
+}
+
+void no_deref_modf() {
+ modf_result_f32 a = modf_result_f32{.fract=0.5f, .whole=1.0f};
+ float fract = a.fract;
+ float whole = a.whole;
+}
+
+void deref_frexp() {
+ frexp_result_f32 a = frexp_result_f32{.fract=0.75f, .exp=1};
+ float fract = a.fract;
+ int exp = a.exp;
+}
+
+void no_deref_frexp() {
+ frexp_result_f32 a = frexp_result_f32{.fract=0.75f, .exp=1};
+ float fract = a.fract;
+ int exp = a.exp;
+}
+
+kernel void tint_symbol() {
+ deref_modf();
+ no_deref_modf();
+ deref_frexp();
+ no_deref_frexp();
+ return;
+}
+
diff --git a/test/tint/ptr_sugar/builtin_struct.wgsl.expected.spvasm b/test/tint/ptr_sugar/builtin_struct.wgsl.expected.spvasm
new file mode 100644
index 0000000..5459222
--- /dev/null
+++ b/test/tint/ptr_sugar/builtin_struct.wgsl.expected.spvasm
@@ -0,0 +1,123 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 66
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %deref_modf "deref_modf"
+ OpName %__modf_result_f32 "__modf_result_f32"
+ OpMemberName %__modf_result_f32 0 "fract"
+ OpMemberName %__modf_result_f32 1 "whole"
+ OpName %a "a"
+ OpName %fract "fract"
+ OpName %whole "whole"
+ OpName %no_deref_modf "no_deref_modf"
+ OpName %a_0 "a"
+ OpName %fract_0 "fract"
+ OpName %whole_0 "whole"
+ OpName %deref_frexp "deref_frexp"
+ OpName %__frexp_result_f32 "__frexp_result_f32"
+ OpMemberName %__frexp_result_f32 0 "fract"
+ OpMemberName %__frexp_result_f32 1 "exp"
+ OpName %a_1 "a"
+ OpName %fract_1 "fract"
+ OpName %exp "exp"
+ OpName %no_deref_frexp "no_deref_frexp"
+ OpName %a_2 "a"
+ OpName %fract_2 "fract"
+ OpName %exp_0 "exp"
+ OpName %main "main"
+ OpMemberDecorate %__modf_result_f32 0 Offset 0
+ OpMemberDecorate %__modf_result_f32 1 Offset 4
+ OpMemberDecorate %__frexp_result_f32 0 Offset 0
+ OpMemberDecorate %__frexp_result_f32 1 Offset 4
+ %void = OpTypeVoid
+ %1 = OpTypeFunction %void
+ %float = OpTypeFloat 32
+%__modf_result_f32 = OpTypeStruct %float %float
+ %float_0_5 = OpConstant %float 0.5
+ %float_1 = OpConstant %float 1
+ %9 = OpConstantComposite %__modf_result_f32 %float_0_5 %float_1
+%_ptr_Function___modf_result_f32 = OpTypePointer Function %__modf_result_f32
+ %12 = OpConstantNull %__modf_result_f32
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+%_ptr_Function_float = OpTypePointer Function %float
+ %19 = OpConstantNull %float
+ %uint_1 = OpConstant %uint 1
+ %int = OpTypeInt 32 1
+%__frexp_result_f32 = OpTypeStruct %float %int
+ %float_0_75 = OpConstant %float 0.75
+ %int_1 = OpConstant %int 1
+ %39 = OpConstantComposite %__frexp_result_f32 %float_0_75 %int_1
+%_ptr_Function___frexp_result_f32 = OpTypePointer Function %__frexp_result_f32
+ %42 = OpConstantNull %__frexp_result_f32
+%_ptr_Function_int = OpTypePointer Function %int
+ %50 = OpConstantNull %int
+ %deref_modf = OpFunction %void None %1
+ %4 = OpLabel
+ %a = OpVariable %_ptr_Function___modf_result_f32 Function %12
+ %fract = OpVariable %_ptr_Function_float Function %19
+ %whole = OpVariable %_ptr_Function_float Function %19
+ OpStore %a %9
+ %16 = OpAccessChain %_ptr_Function_float %a %uint_0
+ %17 = OpLoad %float %16
+ OpStore %fract %17
+ %21 = OpAccessChain %_ptr_Function_float %a %uint_1
+ %22 = OpLoad %float %21
+ OpStore %whole %22
+ OpReturn
+ OpFunctionEnd
+%no_deref_modf = OpFunction %void None %1
+ %25 = OpLabel
+ %a_0 = OpVariable %_ptr_Function___modf_result_f32 Function %12
+ %fract_0 = OpVariable %_ptr_Function_float Function %19
+ %whole_0 = OpVariable %_ptr_Function_float Function %19
+ OpStore %a_0 %9
+ %27 = OpAccessChain %_ptr_Function_float %a_0 %uint_0
+ %28 = OpLoad %float %27
+ OpStore %fract_0 %28
+ %30 = OpAccessChain %_ptr_Function_float %a_0 %uint_1
+ %31 = OpLoad %float %30
+ OpStore %whole_0 %31
+ OpReturn
+ OpFunctionEnd
+%deref_frexp = OpFunction %void None %1
+ %34 = OpLabel
+ %a_1 = OpVariable %_ptr_Function___frexp_result_f32 Function %42
+ %fract_1 = OpVariable %_ptr_Function_float Function %19
+ %exp = OpVariable %_ptr_Function_int Function %50
+ OpStore %a_1 %39
+ %43 = OpAccessChain %_ptr_Function_float %a_1 %uint_0
+ %44 = OpLoad %float %43
+ OpStore %fract_1 %44
+ %47 = OpAccessChain %_ptr_Function_int %a_1 %uint_1
+ %48 = OpLoad %int %47
+ OpStore %exp %48
+ OpReturn
+ OpFunctionEnd
+%no_deref_frexp = OpFunction %void None %1
+ %52 = OpLabel
+ %a_2 = OpVariable %_ptr_Function___frexp_result_f32 Function %42
+ %fract_2 = OpVariable %_ptr_Function_float Function %19
+ %exp_0 = OpVariable %_ptr_Function_int Function %50
+ OpStore %a_2 %39
+ %54 = OpAccessChain %_ptr_Function_float %a_2 %uint_0
+ %55 = OpLoad %float %54
+ OpStore %fract_2 %55
+ %57 = OpAccessChain %_ptr_Function_int %a_2 %uint_1
+ %58 = OpLoad %int %57
+ OpStore %exp_0 %58
+ OpReturn
+ OpFunctionEnd
+ %main = OpFunction %void None %1
+ %61 = OpLabel
+ %62 = OpFunctionCall %void %deref_modf
+ %63 = OpFunctionCall %void %no_deref_modf
+ %64 = OpFunctionCall %void %deref_frexp
+ %65 = OpFunctionCall %void %no_deref_frexp
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/ptr_sugar/builtin_struct.wgsl.expected.wgsl b/test/tint/ptr_sugar/builtin_struct.wgsl.expected.wgsl
new file mode 100644
index 0000000..7349c16
--- /dev/null
+++ b/test/tint/ptr_sugar/builtin_struct.wgsl.expected.wgsl
@@ -0,0 +1,35 @@
+fn deref_modf() {
+ var a = modf(1.5);
+ let p = &(a);
+ var fract = (*(p)).fract;
+ var whole = (*(p)).whole;
+}
+
+fn no_deref_modf() {
+ var a = modf(1.5);
+ let p = &(a);
+ var fract = p.fract;
+ var whole = p.whole;
+}
+
+fn deref_frexp() {
+ var a = frexp(1.5);
+ let p = &(a);
+ var fract = (*(p)).fract;
+ var exp = (*(p)).exp;
+}
+
+fn no_deref_frexp() {
+ var a = frexp(1.5);
+ let p = &(a);
+ var fract = p.fract;
+ var exp = p.exp;
+}
+
+@compute @workgroup_size(1)
+fn main() {
+ deref_modf();
+ no_deref_modf();
+ deref_frexp();
+ no_deref_frexp();
+}
diff --git a/test/tint/ptr_sugar/compound_assign_index.wgsl b/test/tint/ptr_sugar/compound_assign_index.wgsl
new file mode 100644
index 0000000..ff01d4c
--- /dev/null
+++ b/test/tint/ptr_sugar/compound_assign_index.wgsl
@@ -0,0 +1,31 @@
+fn deref() {
+ var a : vec3<i32>;
+ let p = &a;
+ (*p)[0] += 42;
+}
+
+fn no_deref() {
+ var a : vec3<i32>;
+ let p = &a;
+ p[0] += 42;
+}
+
+fn deref_inc() {
+ var a : vec3<i32>;
+ let p = &a;
+ (*p)[0]++;
+}
+
+fn no_deref_inc() {
+ var a : vec3<i32>;
+ let p = &a;
+ p[0]++;
+}
+
+@compute @workgroup_size(1)
+fn main() {
+ deref();
+ no_deref();
+ deref_inc();
+ no_deref_inc();
+}
diff --git a/test/tint/ptr_sugar/compound_assign_index.wgsl.expected.dxc.hlsl b/test/tint/ptr_sugar/compound_assign_index.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..25853e3
--- /dev/null
+++ b/test/tint/ptr_sugar/compound_assign_index.wgsl.expected.dxc.hlsl
@@ -0,0 +1,36 @@
+void set_int3(inout int3 vec, int idx, int val) {
+ vec = (idx.xxx == int3(0, 1, 2)) ? val.xxx : vec;
+}
+
+void deref() {
+ int3 a = int3(0, 0, 0);
+ const int tint_symbol_1 = 0;
+ set_int3(a, tint_symbol_1, (a[tint_symbol_1] + 42));
+}
+
+void no_deref() {
+ int3 a = int3(0, 0, 0);
+ const int tint_symbol_3 = 0;
+ set_int3(a, tint_symbol_3, (a[tint_symbol_3] + 42));
+}
+
+void deref_inc() {
+ int3 a = int3(0, 0, 0);
+ const int tint_symbol_5 = 0;
+ set_int3(a, tint_symbol_5, (a[tint_symbol_5] + 1));
+}
+
+void no_deref_inc() {
+ int3 a = int3(0, 0, 0);
+ const int tint_symbol_7 = 0;
+ set_int3(a, tint_symbol_7, (a[tint_symbol_7] + 1));
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+ deref();
+ no_deref();
+ deref_inc();
+ no_deref_inc();
+ return;
+}
diff --git a/test/tint/ptr_sugar/compound_assign_index.wgsl.expected.fxc.hlsl b/test/tint/ptr_sugar/compound_assign_index.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..25853e3
--- /dev/null
+++ b/test/tint/ptr_sugar/compound_assign_index.wgsl.expected.fxc.hlsl
@@ -0,0 +1,36 @@
+void set_int3(inout int3 vec, int idx, int val) {
+ vec = (idx.xxx == int3(0, 1, 2)) ? val.xxx : vec;
+}
+
+void deref() {
+ int3 a = int3(0, 0, 0);
+ const int tint_symbol_1 = 0;
+ set_int3(a, tint_symbol_1, (a[tint_symbol_1] + 42));
+}
+
+void no_deref() {
+ int3 a = int3(0, 0, 0);
+ const int tint_symbol_3 = 0;
+ set_int3(a, tint_symbol_3, (a[tint_symbol_3] + 42));
+}
+
+void deref_inc() {
+ int3 a = int3(0, 0, 0);
+ const int tint_symbol_5 = 0;
+ set_int3(a, tint_symbol_5, (a[tint_symbol_5] + 1));
+}
+
+void no_deref_inc() {
+ int3 a = int3(0, 0, 0);
+ const int tint_symbol_7 = 0;
+ set_int3(a, tint_symbol_7, (a[tint_symbol_7] + 1));
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+ deref();
+ no_deref();
+ deref_inc();
+ no_deref_inc();
+ return;
+}
diff --git a/test/tint/ptr_sugar/compound_assign_index.wgsl.expected.glsl b/test/tint/ptr_sugar/compound_assign_index.wgsl.expected.glsl
new file mode 100644
index 0000000..dd04e0f
--- /dev/null
+++ b/test/tint/ptr_sugar/compound_assign_index.wgsl.expected.glsl
@@ -0,0 +1,38 @@
+#version 310 es
+
+void deref() {
+ ivec3 a = ivec3(0, 0, 0);
+ int tint_symbol_2 = 0;
+ a[tint_symbol_2] = (a[tint_symbol_2] + 42);
+}
+
+void no_deref() {
+ ivec3 a = ivec3(0, 0, 0);
+ int tint_symbol_4 = 0;
+ a[tint_symbol_4] = (a[tint_symbol_4] + 42);
+}
+
+void deref_inc() {
+ ivec3 a = ivec3(0, 0, 0);
+ int tint_symbol_6 = 0;
+ a[tint_symbol_6] = (a[tint_symbol_6] + 1);
+}
+
+void no_deref_inc() {
+ ivec3 a = ivec3(0, 0, 0);
+ int tint_symbol_8 = 0;
+ a[tint_symbol_8] = (a[tint_symbol_8] + 1);
+}
+
+void tint_symbol() {
+ deref();
+ no_deref();
+ deref_inc();
+ no_deref_inc();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ tint_symbol();
+ return;
+}
diff --git a/test/tint/ptr_sugar/compound_assign_index.wgsl.expected.msl b/test/tint/ptr_sugar/compound_assign_index.wgsl.expected.msl
new file mode 100644
index 0000000..0a18439
--- /dev/null
+++ b/test/tint/ptr_sugar/compound_assign_index.wgsl.expected.msl
@@ -0,0 +1,35 @@
+#include <metal_stdlib>
+
+using namespace metal;
+void deref() {
+ int3 a = 0;
+ int const tint_symbol_2 = 0;
+ a[tint_symbol_2] = as_type<int>((as_type<uint>(a[tint_symbol_2]) + as_type<uint>(42)));
+}
+
+void no_deref() {
+ int3 a = 0;
+ int const tint_symbol_4 = 0;
+ a[tint_symbol_4] = as_type<int>((as_type<uint>(a[tint_symbol_4]) + as_type<uint>(42)));
+}
+
+void deref_inc() {
+ int3 a = 0;
+ int const tint_symbol_6 = 0;
+ a[tint_symbol_6] = as_type<int>((as_type<uint>(a[tint_symbol_6]) + as_type<uint>(1)));
+}
+
+void no_deref_inc() {
+ int3 a = 0;
+ int const tint_symbol_8 = 0;
+ a[tint_symbol_8] = as_type<int>((as_type<uint>(a[tint_symbol_8]) + as_type<uint>(1)));
+}
+
+kernel void tint_symbol() {
+ deref();
+ no_deref();
+ deref_inc();
+ no_deref_inc();
+ return;
+}
+
diff --git a/test/tint/ptr_sugar/compound_assign_index.wgsl.expected.spvasm b/test/tint/ptr_sugar/compound_assign_index.wgsl.expected.spvasm
new file mode 100644
index 0000000..237c95a
--- /dev/null
+++ b/test/tint/ptr_sugar/compound_assign_index.wgsl.expected.spvasm
@@ -0,0 +1,76 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 45
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %deref "deref"
+ OpName %a "a"
+ OpName %no_deref "no_deref"
+ OpName %a_0 "a"
+ OpName %deref_inc "deref_inc"
+ OpName %a_1 "a"
+ OpName %no_deref_inc "no_deref_inc"
+ OpName %a_2 "a"
+ OpName %main "main"
+ %void = OpTypeVoid
+ %1 = OpTypeFunction %void
+ %int = OpTypeInt 32 1
+ %v3int = OpTypeVector %int 3
+%_ptr_Function_v3int = OpTypePointer Function %v3int
+ %9 = OpConstantNull %v3int
+ %10 = OpConstantNull %int
+%_ptr_Function_int = OpTypePointer Function %int
+ %int_42 = OpConstant %int 42
+ %int_1 = OpConstant %int 1
+ %deref = OpFunction %void None %1
+ %4 = OpLabel
+ %a = OpVariable %_ptr_Function_v3int Function %9
+ %12 = OpAccessChain %_ptr_Function_int %a %10
+ %13 = OpAccessChain %_ptr_Function_int %a %10
+ %14 = OpLoad %int %13
+ %16 = OpIAdd %int %14 %int_42
+ OpStore %12 %16
+ OpReturn
+ OpFunctionEnd
+ %no_deref = OpFunction %void None %1
+ %18 = OpLabel
+ %a_0 = OpVariable %_ptr_Function_v3int Function %9
+ %20 = OpAccessChain %_ptr_Function_int %a_0 %10
+ %21 = OpAccessChain %_ptr_Function_int %a_0 %10
+ %22 = OpLoad %int %21
+ %23 = OpIAdd %int %22 %int_42
+ OpStore %20 %23
+ OpReturn
+ OpFunctionEnd
+ %deref_inc = OpFunction %void None %1
+ %25 = OpLabel
+ %a_1 = OpVariable %_ptr_Function_v3int Function %9
+ %27 = OpAccessChain %_ptr_Function_int %a_1 %10
+ %28 = OpAccessChain %_ptr_Function_int %a_1 %10
+ %29 = OpLoad %int %28
+ %31 = OpIAdd %int %29 %int_1
+ OpStore %27 %31
+ OpReturn
+ OpFunctionEnd
+%no_deref_inc = OpFunction %void None %1
+ %33 = OpLabel
+ %a_2 = OpVariable %_ptr_Function_v3int Function %9
+ %35 = OpAccessChain %_ptr_Function_int %a_2 %10
+ %36 = OpAccessChain %_ptr_Function_int %a_2 %10
+ %37 = OpLoad %int %36
+ %38 = OpIAdd %int %37 %int_1
+ OpStore %35 %38
+ OpReturn
+ OpFunctionEnd
+ %main = OpFunction %void None %1
+ %40 = OpLabel
+ %41 = OpFunctionCall %void %deref
+ %42 = OpFunctionCall %void %no_deref
+ %43 = OpFunctionCall %void %deref_inc
+ %44 = OpFunctionCall %void %no_deref_inc
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/ptr_sugar/compound_assign_index.wgsl.expected.wgsl b/test/tint/ptr_sugar/compound_assign_index.wgsl.expected.wgsl
new file mode 100644
index 0000000..9fad9aa
--- /dev/null
+++ b/test/tint/ptr_sugar/compound_assign_index.wgsl.expected.wgsl
@@ -0,0 +1,31 @@
+fn deref() {
+ var a : vec3<i32>;
+ let p = &(a);
+ (*(p))[0] += 42;
+}
+
+fn no_deref() {
+ var a : vec3<i32>;
+ let p = &(a);
+ p[0] += 42;
+}
+
+fn deref_inc() {
+ var a : vec3<i32>;
+ let p = &(a);
+ (*(p))[0]++;
+}
+
+fn no_deref_inc() {
+ var a : vec3<i32>;
+ let p = &(a);
+ p[0]++;
+}
+
+@compute @workgroup_size(1)
+fn main() {
+ deref();
+ no_deref();
+ deref_inc();
+ no_deref_inc();
+}
diff --git a/test/tint/ptr_sugar/compound_assign_member.wgsl b/test/tint/ptr_sugar/compound_assign_member.wgsl
new file mode 100644
index 0000000..2785bd5
--- /dev/null
+++ b/test/tint/ptr_sugar/compound_assign_member.wgsl
@@ -0,0 +1,17 @@
+fn deref() {
+ var a : vec3<i32>;
+ let p = &a;
+ (*p).x += 42;
+}
+
+fn no_deref() {
+ var a : vec3<i32>;
+ let p = &a;
+ p.x += 42;
+}
+
+@compute @workgroup_size(1)
+fn main() {
+ deref();
+ no_deref();
+}
diff --git a/test/tint/ptr_sugar/compound_assign_member.wgsl.expected.dxc.hlsl b/test/tint/ptr_sugar/compound_assign_member.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..630d44f
--- /dev/null
+++ b/test/tint/ptr_sugar/compound_assign_member.wgsl.expected.dxc.hlsl
@@ -0,0 +1,16 @@
+void deref() {
+ int3 a = int3(0, 0, 0);
+ a.x = (a.x + 42);
+}
+
+void no_deref() {
+ int3 a = int3(0, 0, 0);
+ a.x = (a.x + 42);
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+ deref();
+ no_deref();
+ return;
+}
diff --git a/test/tint/ptr_sugar/compound_assign_member.wgsl.expected.fxc.hlsl b/test/tint/ptr_sugar/compound_assign_member.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..630d44f
--- /dev/null
+++ b/test/tint/ptr_sugar/compound_assign_member.wgsl.expected.fxc.hlsl
@@ -0,0 +1,16 @@
+void deref() {
+ int3 a = int3(0, 0, 0);
+ a.x = (a.x + 42);
+}
+
+void no_deref() {
+ int3 a = int3(0, 0, 0);
+ a.x = (a.x + 42);
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+ deref();
+ no_deref();
+ return;
+}
diff --git a/test/tint/ptr_sugar/compound_assign_member.wgsl.expected.glsl b/test/tint/ptr_sugar/compound_assign_member.wgsl.expected.glsl
new file mode 100644
index 0000000..62a3b1f
--- /dev/null
+++ b/test/tint/ptr_sugar/compound_assign_member.wgsl.expected.glsl
@@ -0,0 +1,22 @@
+#version 310 es
+
+void deref() {
+ ivec3 a = ivec3(0, 0, 0);
+ a.x = (a.x + 42);
+}
+
+void no_deref() {
+ ivec3 a = ivec3(0, 0, 0);
+ a.x = (a.x + 42);
+}
+
+void tint_symbol() {
+ deref();
+ no_deref();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ tint_symbol();
+ return;
+}
diff --git a/test/tint/ptr_sugar/compound_assign_member.wgsl.expected.msl b/test/tint/ptr_sugar/compound_assign_member.wgsl.expected.msl
new file mode 100644
index 0000000..65e3c9e
--- /dev/null
+++ b/test/tint/ptr_sugar/compound_assign_member.wgsl.expected.msl
@@ -0,0 +1,19 @@
+#include <metal_stdlib>
+
+using namespace metal;
+void deref() {
+ int3 a = 0;
+ a[0] = as_type<int>((as_type<uint>(a[0]) + as_type<uint>(42)));
+}
+
+void no_deref() {
+ int3 a = 0;
+ a[0] = as_type<int>((as_type<uint>(a[0]) + as_type<uint>(42)));
+}
+
+kernel void tint_symbol() {
+ deref();
+ no_deref();
+ return;
+}
+
diff --git a/test/tint/ptr_sugar/compound_assign_member.wgsl.expected.spvasm b/test/tint/ptr_sugar/compound_assign_member.wgsl.expected.spvasm
new file mode 100644
index 0000000..7e3b8ff
--- /dev/null
+++ b/test/tint/ptr_sugar/compound_assign_member.wgsl.expected.spvasm
@@ -0,0 +1,50 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 29
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %deref "deref"
+ OpName %a "a"
+ OpName %no_deref "no_deref"
+ OpName %a_0 "a"
+ OpName %main "main"
+ %void = OpTypeVoid
+ %1 = OpTypeFunction %void
+ %int = OpTypeInt 32 1
+ %v3int = OpTypeVector %int 3
+%_ptr_Function_v3int = OpTypePointer Function %v3int
+ %9 = OpConstantNull %v3int
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+%_ptr_Function_int = OpTypePointer Function %int
+ %int_42 = OpConstant %int 42
+ %deref = OpFunction %void None %1
+ %4 = OpLabel
+ %a = OpVariable %_ptr_Function_v3int Function %9
+ %13 = OpAccessChain %_ptr_Function_int %a %uint_0
+ %14 = OpAccessChain %_ptr_Function_int %a %uint_0
+ %15 = OpLoad %int %14
+ %17 = OpIAdd %int %15 %int_42
+ OpStore %13 %17
+ OpReturn
+ OpFunctionEnd
+ %no_deref = OpFunction %void None %1
+ %19 = OpLabel
+ %a_0 = OpVariable %_ptr_Function_v3int Function %9
+ %21 = OpAccessChain %_ptr_Function_int %a_0 %uint_0
+ %22 = OpAccessChain %_ptr_Function_int %a_0 %uint_0
+ %23 = OpLoad %int %22
+ %24 = OpIAdd %int %23 %int_42
+ OpStore %21 %24
+ OpReturn
+ OpFunctionEnd
+ %main = OpFunction %void None %1
+ %26 = OpLabel
+ %27 = OpFunctionCall %void %deref
+ %28 = OpFunctionCall %void %no_deref
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/ptr_sugar/compound_assign_member.wgsl.expected.wgsl b/test/tint/ptr_sugar/compound_assign_member.wgsl.expected.wgsl
new file mode 100644
index 0000000..3a01d23
--- /dev/null
+++ b/test/tint/ptr_sugar/compound_assign_member.wgsl.expected.wgsl
@@ -0,0 +1,17 @@
+fn deref() {
+ var a : vec3<i32>;
+ let p = &(a);
+ (*(p)).x += 42;
+}
+
+fn no_deref() {
+ var a : vec3<i32>;
+ let p = &(a);
+ p.x += 42;
+}
+
+@compute @workgroup_size(1)
+fn main() {
+ deref();
+ no_deref();
+}
diff --git a/test/tint/ptr_sugar/matrix.wgsl b/test/tint/ptr_sugar/matrix.wgsl
new file mode 100644
index 0000000..5a768d9
--- /dev/null
+++ b/test/tint/ptr_sugar/matrix.wgsl
@@ -0,0 +1,19 @@
+fn deref() {
+ var a : mat2x3<f32>;
+ let p = &a;
+ var b = (*p)[0];
+ (*p)[0] = vec3<f32>(1.0, 2.0, 3.0);
+}
+
+fn no_deref() {
+ var a : mat2x3<f32>;
+ let p = &a;
+ var b = p[0];
+ p[0] = vec3<f32>(1.0, 2.0, 3.0);
+}
+
+@compute @workgroup_size(1)
+fn main() {
+ deref();
+ no_deref();
+}
diff --git a/test/tint/ptr_sugar/matrix.wgsl.expected.dxc.hlsl b/test/tint/ptr_sugar/matrix.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..7790b20
--- /dev/null
+++ b/test/tint/ptr_sugar/matrix.wgsl.expected.dxc.hlsl
@@ -0,0 +1,18 @@
+void deref() {
+ float2x3 a = float2x3(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f);
+ float3 b = a[0];
+ a[0] = float3(1.0f, 2.0f, 3.0f);
+}
+
+void no_deref() {
+ float2x3 a = float2x3(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f);
+ float3 b = a[0];
+ a[0] = float3(1.0f, 2.0f, 3.0f);
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+ deref();
+ no_deref();
+ return;
+}
diff --git a/test/tint/ptr_sugar/matrix.wgsl.expected.fxc.hlsl b/test/tint/ptr_sugar/matrix.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..7790b20
--- /dev/null
+++ b/test/tint/ptr_sugar/matrix.wgsl.expected.fxc.hlsl
@@ -0,0 +1,18 @@
+void deref() {
+ float2x3 a = float2x3(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f);
+ float3 b = a[0];
+ a[0] = float3(1.0f, 2.0f, 3.0f);
+}
+
+void no_deref() {
+ float2x3 a = float2x3(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f);
+ float3 b = a[0];
+ a[0] = float3(1.0f, 2.0f, 3.0f);
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+ deref();
+ no_deref();
+ return;
+}
diff --git a/test/tint/ptr_sugar/matrix.wgsl.expected.glsl b/test/tint/ptr_sugar/matrix.wgsl.expected.glsl
new file mode 100644
index 0000000..25cae38
--- /dev/null
+++ b/test/tint/ptr_sugar/matrix.wgsl.expected.glsl
@@ -0,0 +1,24 @@
+#version 310 es
+
+void deref() {
+ mat2x3 a = mat2x3(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f);
+ vec3 b = a[0];
+ a[0] = vec3(1.0f, 2.0f, 3.0f);
+}
+
+void no_deref() {
+ mat2x3 a = mat2x3(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f);
+ vec3 b = a[0];
+ a[0] = vec3(1.0f, 2.0f, 3.0f);
+}
+
+void tint_symbol() {
+ deref();
+ no_deref();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ tint_symbol();
+ return;
+}
diff --git a/test/tint/ptr_sugar/matrix.wgsl.expected.msl b/test/tint/ptr_sugar/matrix.wgsl.expected.msl
new file mode 100644
index 0000000..65ab860
--- /dev/null
+++ b/test/tint/ptr_sugar/matrix.wgsl.expected.msl
@@ -0,0 +1,21 @@
+#include <metal_stdlib>
+
+using namespace metal;
+void deref() {
+ float2x3 a = float2x3(0.0f);
+ float3 b = a[0];
+ a[0] = float3(1.0f, 2.0f, 3.0f);
+}
+
+void no_deref() {
+ float2x3 a = float2x3(0.0f);
+ float3 b = a[0];
+ a[0] = float3(1.0f, 2.0f, 3.0f);
+}
+
+kernel void tint_symbol() {
+ deref();
+ no_deref();
+ return;
+}
+
diff --git a/test/tint/ptr_sugar/matrix.wgsl.expected.spvasm b/test/tint/ptr_sugar/matrix.wgsl.expected.spvasm
new file mode 100644
index 0000000..4f1e62d
--- /dev/null
+++ b/test/tint/ptr_sugar/matrix.wgsl.expected.spvasm
@@ -0,0 +1,59 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 34
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %deref "deref"
+ OpName %a "a"
+ OpName %b "b"
+ OpName %no_deref "no_deref"
+ OpName %a_0 "a"
+ OpName %b_0 "b"
+ OpName %main "main"
+ %void = OpTypeVoid
+ %1 = OpTypeFunction %void
+ %float = OpTypeFloat 32
+ %v3float = OpTypeVector %float 3
+%mat2v3float = OpTypeMatrix %v3float 2
+%_ptr_Function_mat2v3float = OpTypePointer Function %mat2v3float
+ %10 = OpConstantNull %mat2v3float
+ %int = OpTypeInt 32 1
+ %12 = OpConstantNull %int
+%_ptr_Function_v3float = OpTypePointer Function %v3float
+ %17 = OpConstantNull %v3float
+ %float_1 = OpConstant %float 1
+ %float_2 = OpConstant %float 2
+ %float_3 = OpConstant %float 3
+ %22 = OpConstantComposite %v3float %float_1 %float_2 %float_3
+ %deref = OpFunction %void None %1
+ %4 = OpLabel
+ %a = OpVariable %_ptr_Function_mat2v3float Function %10
+ %b = OpVariable %_ptr_Function_v3float Function %17
+ %14 = OpAccessChain %_ptr_Function_v3float %a %12
+ %15 = OpLoad %v3float %14
+ OpStore %b %15
+ %18 = OpAccessChain %_ptr_Function_v3float %a %12
+ OpStore %18 %22
+ OpReturn
+ OpFunctionEnd
+ %no_deref = OpFunction %void None %1
+ %24 = OpLabel
+ %a_0 = OpVariable %_ptr_Function_mat2v3float Function %10
+ %b_0 = OpVariable %_ptr_Function_v3float Function %17
+ %26 = OpAccessChain %_ptr_Function_v3float %a_0 %12
+ %27 = OpLoad %v3float %26
+ OpStore %b_0 %27
+ %29 = OpAccessChain %_ptr_Function_v3float %a_0 %12
+ OpStore %29 %22
+ OpReturn
+ OpFunctionEnd
+ %main = OpFunction %void None %1
+ %31 = OpLabel
+ %32 = OpFunctionCall %void %deref
+ %33 = OpFunctionCall %void %no_deref
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/ptr_sugar/matrix.wgsl.expected.wgsl b/test/tint/ptr_sugar/matrix.wgsl.expected.wgsl
new file mode 100644
index 0000000..3b18d24
--- /dev/null
+++ b/test/tint/ptr_sugar/matrix.wgsl.expected.wgsl
@@ -0,0 +1,19 @@
+fn deref() {
+ var a : mat2x3<f32>;
+ let p = &(a);
+ var b = (*(p))[0];
+ (*(p))[0] = vec3<f32>(1.0, 2.0, 3.0);
+}
+
+fn no_deref() {
+ var a : mat2x3<f32>;
+ let p = &(a);
+ var b = p[0];
+ p[0] = vec3<f32>(1.0, 2.0, 3.0);
+}
+
+@compute @workgroup_size(1)
+fn main() {
+ deref();
+ no_deref();
+}
diff --git a/test/tint/ptr_sugar/struct.wgsl b/test/tint/ptr_sugar/struct.wgsl
new file mode 100644
index 0000000..1c6c7fe
--- /dev/null
+++ b/test/tint/ptr_sugar/struct.wgsl
@@ -0,0 +1,23 @@
+struct S {
+ x : i32,
+}
+
+fn deref() {
+ var a : S;
+ let p = &a;
+ var b = (*p).x;
+ (*p).x = 42;
+}
+
+fn no_deref() {
+ var a : S;
+ let p = &a;
+ var b = p.x;
+ p.x = 42;
+}
+
+@compute @workgroup_size(1)
+fn main() {
+ deref();
+ no_deref();
+}
diff --git a/test/tint/ptr_sugar/struct.wgsl.expected.dxc.hlsl b/test/tint/ptr_sugar/struct.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..628808f
--- /dev/null
+++ b/test/tint/ptr_sugar/struct.wgsl.expected.dxc.hlsl
@@ -0,0 +1,22 @@
+struct S {
+ int x;
+};
+
+void deref() {
+ S a = (S)0;
+ int b = a.x;
+ a.x = 42;
+}
+
+void no_deref() {
+ S a = (S)0;
+ int b = a.x;
+ a.x = 42;
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+ deref();
+ no_deref();
+ return;
+}
diff --git a/test/tint/ptr_sugar/struct.wgsl.expected.fxc.hlsl b/test/tint/ptr_sugar/struct.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..628808f
--- /dev/null
+++ b/test/tint/ptr_sugar/struct.wgsl.expected.fxc.hlsl
@@ -0,0 +1,22 @@
+struct S {
+ int x;
+};
+
+void deref() {
+ S a = (S)0;
+ int b = a.x;
+ a.x = 42;
+}
+
+void no_deref() {
+ S a = (S)0;
+ int b = a.x;
+ a.x = 42;
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+ deref();
+ no_deref();
+ return;
+}
diff --git a/test/tint/ptr_sugar/struct.wgsl.expected.glsl b/test/tint/ptr_sugar/struct.wgsl.expected.glsl
new file mode 100644
index 0000000..dceeb2c
--- /dev/null
+++ b/test/tint/ptr_sugar/struct.wgsl.expected.glsl
@@ -0,0 +1,28 @@
+#version 310 es
+
+struct S {
+ int x;
+};
+
+void deref() {
+ S a = S(0);
+ int b = a.x;
+ a.x = 42;
+}
+
+void no_deref() {
+ S a = S(0);
+ int b = a.x;
+ a.x = 42;
+}
+
+void tint_symbol() {
+ deref();
+ no_deref();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ tint_symbol();
+ return;
+}
diff --git a/test/tint/ptr_sugar/struct.wgsl.expected.msl b/test/tint/ptr_sugar/struct.wgsl.expected.msl
new file mode 100644
index 0000000..1b6ebcd
--- /dev/null
+++ b/test/tint/ptr_sugar/struct.wgsl.expected.msl
@@ -0,0 +1,25 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct S {
+ int x;
+};
+
+void deref() {
+ S a = {};
+ int b = a.x;
+ a.x = 42;
+}
+
+void no_deref() {
+ S a = {};
+ int b = a.x;
+ a.x = 42;
+}
+
+kernel void tint_symbol() {
+ deref();
+ no_deref();
+ return;
+}
+
diff --git a/test/tint/ptr_sugar/struct.wgsl.expected.spvasm b/test/tint/ptr_sugar/struct.wgsl.expected.spvasm
new file mode 100644
index 0000000..d4c38c9
--- /dev/null
+++ b/test/tint/ptr_sugar/struct.wgsl.expected.spvasm
@@ -0,0 +1,58 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 30
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %deref "deref"
+ OpName %S "S"
+ OpMemberName %S 0 "x"
+ OpName %a "a"
+ OpName %b "b"
+ OpName %no_deref "no_deref"
+ OpName %a_0 "a"
+ OpName %b_0 "b"
+ OpName %main "main"
+ OpMemberDecorate %S 0 Offset 0
+ %void = OpTypeVoid
+ %1 = OpTypeFunction %void
+ %int = OpTypeInt 32 1
+ %S = OpTypeStruct %int
+%_ptr_Function_S = OpTypePointer Function %S
+ %9 = OpConstantNull %S
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+%_ptr_Function_int = OpTypePointer Function %int
+ %16 = OpConstantNull %int
+ %int_42 = OpConstant %int 42
+ %deref = OpFunction %void None %1
+ %4 = OpLabel
+ %a = OpVariable %_ptr_Function_S Function %9
+ %b = OpVariable %_ptr_Function_int Function %16
+ %13 = OpAccessChain %_ptr_Function_int %a %uint_0
+ %14 = OpLoad %int %13
+ OpStore %b %14
+ %17 = OpAccessChain %_ptr_Function_int %a %uint_0
+ OpStore %17 %int_42
+ OpReturn
+ OpFunctionEnd
+ %no_deref = OpFunction %void None %1
+ %20 = OpLabel
+ %a_0 = OpVariable %_ptr_Function_S Function %9
+ %b_0 = OpVariable %_ptr_Function_int Function %16
+ %22 = OpAccessChain %_ptr_Function_int %a_0 %uint_0
+ %23 = OpLoad %int %22
+ OpStore %b_0 %23
+ %25 = OpAccessChain %_ptr_Function_int %a_0 %uint_0
+ OpStore %25 %int_42
+ OpReturn
+ OpFunctionEnd
+ %main = OpFunction %void None %1
+ %27 = OpLabel
+ %28 = OpFunctionCall %void %deref
+ %29 = OpFunctionCall %void %no_deref
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/ptr_sugar/struct.wgsl.expected.wgsl b/test/tint/ptr_sugar/struct.wgsl.expected.wgsl
new file mode 100644
index 0000000..e01aecd
--- /dev/null
+++ b/test/tint/ptr_sugar/struct.wgsl.expected.wgsl
@@ -0,0 +1,23 @@
+struct S {
+ x : i32,
+}
+
+fn deref() {
+ var a : S;
+ let p = &(a);
+ var b = (*(p)).x;
+ (*(p)).x = 42;
+}
+
+fn no_deref() {
+ var a : S;
+ let p = &(a);
+ var b = p.x;
+ p.x = 42;
+}
+
+@compute @workgroup_size(1)
+fn main() {
+ deref();
+ no_deref();
+}
diff --git a/test/tint/ptr_sugar/vector_index.wgsl b/test/tint/ptr_sugar/vector_index.wgsl
new file mode 100644
index 0000000..c4a549a
--- /dev/null
+++ b/test/tint/ptr_sugar/vector_index.wgsl
@@ -0,0 +1,55 @@
+fn deref_const() {
+ var a : vec3<i32>;
+ let p = &a;
+ var b = (*p)[0];
+ (*p)[0] = 42;
+}
+
+fn no_deref_const() {
+ var a : vec3<i32>;
+ let p = &a;
+ var b = p[0];
+ p[0] = 42;
+}
+
+fn deref_let() {
+ var a : vec3<i32>;
+ let p = &a;
+ let i = 0;
+ var b = (*p)[i];
+ (*p)[0] = 42;
+}
+
+fn no_deref_let() {
+ var a : vec3<i32>;
+ let p = &a;
+ let i = 0;
+ var b = p[i];
+ p[0] = 42;
+}
+
+fn deref_var() {
+ var a : vec3<i32>;
+ let p = &a;
+ var i = 0;
+ var b = (*p)[i];
+ (*p)[0] = 42;
+}
+
+fn no_deref_var() {
+ var a : vec3<i32>;
+ let p = &a;
+ let i = 0;
+ var b = p[i];
+ p[0] = 42;
+}
+
+@compute @workgroup_size(1)
+fn main() {
+ deref_const();
+ no_deref_const();
+ deref_let();
+ no_deref_let();
+ deref_var();
+ no_deref_var();
+}
diff --git a/test/tint/ptr_sugar/vector_index.wgsl.expected.dxc.hlsl b/test/tint/ptr_sugar/vector_index.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..7d0de2e
--- /dev/null
+++ b/test/tint/ptr_sugar/vector_index.wgsl.expected.dxc.hlsl
@@ -0,0 +1,50 @@
+void deref_const() {
+ int3 a = int3(0, 0, 0);
+ int b = a[0];
+ a[0] = 42;
+}
+
+void no_deref_const() {
+ int3 a = int3(0, 0, 0);
+ int b = a[0];
+ a[0] = 42;
+}
+
+void deref_let() {
+ int3 a = int3(0, 0, 0);
+ const int i = 0;
+ int b = a[i];
+ a[0] = 42;
+}
+
+void no_deref_let() {
+ int3 a = int3(0, 0, 0);
+ const int i = 0;
+ int b = a[i];
+ a[0] = 42;
+}
+
+void deref_var() {
+ int3 a = int3(0, 0, 0);
+ int i = 0;
+ int b = a[i];
+ a[0] = 42;
+}
+
+void no_deref_var() {
+ int3 a = int3(0, 0, 0);
+ const int i = 0;
+ int b = a[i];
+ a[0] = 42;
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+ deref_const();
+ no_deref_const();
+ deref_let();
+ no_deref_let();
+ deref_var();
+ no_deref_var();
+ return;
+}
diff --git a/test/tint/ptr_sugar/vector_index.wgsl.expected.fxc.hlsl b/test/tint/ptr_sugar/vector_index.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..7d0de2e
--- /dev/null
+++ b/test/tint/ptr_sugar/vector_index.wgsl.expected.fxc.hlsl
@@ -0,0 +1,50 @@
+void deref_const() {
+ int3 a = int3(0, 0, 0);
+ int b = a[0];
+ a[0] = 42;
+}
+
+void no_deref_const() {
+ int3 a = int3(0, 0, 0);
+ int b = a[0];
+ a[0] = 42;
+}
+
+void deref_let() {
+ int3 a = int3(0, 0, 0);
+ const int i = 0;
+ int b = a[i];
+ a[0] = 42;
+}
+
+void no_deref_let() {
+ int3 a = int3(0, 0, 0);
+ const int i = 0;
+ int b = a[i];
+ a[0] = 42;
+}
+
+void deref_var() {
+ int3 a = int3(0, 0, 0);
+ int i = 0;
+ int b = a[i];
+ a[0] = 42;
+}
+
+void no_deref_var() {
+ int3 a = int3(0, 0, 0);
+ const int i = 0;
+ int b = a[i];
+ a[0] = 42;
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+ deref_const();
+ no_deref_const();
+ deref_let();
+ no_deref_let();
+ deref_var();
+ no_deref_var();
+ return;
+}
diff --git a/test/tint/ptr_sugar/vector_index.wgsl.expected.glsl b/test/tint/ptr_sugar/vector_index.wgsl.expected.glsl
new file mode 100644
index 0000000..b78430b
--- /dev/null
+++ b/test/tint/ptr_sugar/vector_index.wgsl.expected.glsl
@@ -0,0 +1,56 @@
+#version 310 es
+
+void deref_const() {
+ ivec3 a = ivec3(0, 0, 0);
+ int b = a[0];
+ a[0] = 42;
+}
+
+void no_deref_const() {
+ ivec3 a = ivec3(0, 0, 0);
+ int b = a[0];
+ a[0] = 42;
+}
+
+void deref_let() {
+ ivec3 a = ivec3(0, 0, 0);
+ int i = 0;
+ int b = a[i];
+ a[0] = 42;
+}
+
+void no_deref_let() {
+ ivec3 a = ivec3(0, 0, 0);
+ int i = 0;
+ int b = a[i];
+ a[0] = 42;
+}
+
+void deref_var() {
+ ivec3 a = ivec3(0, 0, 0);
+ int i = 0;
+ int b = a[i];
+ a[0] = 42;
+}
+
+void no_deref_var() {
+ ivec3 a = ivec3(0, 0, 0);
+ int i = 0;
+ int b = a[i];
+ a[0] = 42;
+}
+
+void tint_symbol() {
+ deref_const();
+ no_deref_const();
+ deref_let();
+ no_deref_let();
+ deref_var();
+ no_deref_var();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ tint_symbol();
+ return;
+}
diff --git a/test/tint/ptr_sugar/vector_index.wgsl.expected.msl b/test/tint/ptr_sugar/vector_index.wgsl.expected.msl
new file mode 100644
index 0000000..2d99f38
--- /dev/null
+++ b/test/tint/ptr_sugar/vector_index.wgsl.expected.msl
@@ -0,0 +1,53 @@
+#include <metal_stdlib>
+
+using namespace metal;
+void deref_const() {
+ int3 a = 0;
+ int b = a[0];
+ a[0] = 42;
+}
+
+void no_deref_const() {
+ int3 a = 0;
+ int b = a[0];
+ a[0] = 42;
+}
+
+void deref_let() {
+ int3 a = 0;
+ int const i = 0;
+ int b = a[i];
+ a[0] = 42;
+}
+
+void no_deref_let() {
+ int3 a = 0;
+ int const i = 0;
+ int b = a[i];
+ a[0] = 42;
+}
+
+void deref_var() {
+ int3 a = 0;
+ int i = 0;
+ int b = a[i];
+ a[0] = 42;
+}
+
+void no_deref_var() {
+ int3 a = 0;
+ int const i = 0;
+ int b = a[i];
+ a[0] = 42;
+}
+
+kernel void tint_symbol() {
+ deref_const();
+ no_deref_const();
+ deref_let();
+ no_deref_let();
+ deref_var();
+ no_deref_var();
+ return;
+}
+
diff --git a/test/tint/ptr_sugar/vector_index.wgsl.expected.spvasm b/test/tint/ptr_sugar/vector_index.wgsl.expected.spvasm
new file mode 100644
index 0000000..7706621
--- /dev/null
+++ b/test/tint/ptr_sugar/vector_index.wgsl.expected.spvasm
@@ -0,0 +1,117 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 62
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %deref_const "deref_const"
+ OpName %a "a"
+ OpName %b "b"
+ OpName %no_deref_const "no_deref_const"
+ OpName %a_0 "a"
+ OpName %b_0 "b"
+ OpName %deref_let "deref_let"
+ OpName %a_1 "a"
+ OpName %b_1 "b"
+ OpName %no_deref_let "no_deref_let"
+ OpName %a_2 "a"
+ OpName %b_2 "b"
+ OpName %deref_var "deref_var"
+ OpName %a_3 "a"
+ OpName %i "i"
+ OpName %b_3 "b"
+ OpName %no_deref_var "no_deref_var"
+ OpName %a_4 "a"
+ OpName %b_4 "b"
+ OpName %main "main"
+ %void = OpTypeVoid
+ %1 = OpTypeFunction %void
+ %int = OpTypeInt 32 1
+ %v3int = OpTypeVector %int 3
+%_ptr_Function_v3int = OpTypePointer Function %v3int
+ %9 = OpConstantNull %v3int
+ %10 = OpConstantNull %int
+%_ptr_Function_int = OpTypePointer Function %int
+ %int_42 = OpConstant %int 42
+%deref_const = OpFunction %void None %1
+ %4 = OpLabel
+ %a = OpVariable %_ptr_Function_v3int Function %9
+ %b = OpVariable %_ptr_Function_int Function %10
+ %12 = OpAccessChain %_ptr_Function_int %a %10
+ %13 = OpLoad %int %12
+ OpStore %b %13
+ %15 = OpAccessChain %_ptr_Function_int %a %10
+ OpStore %15 %int_42
+ OpReturn
+ OpFunctionEnd
+%no_deref_const = OpFunction %void None %1
+ %18 = OpLabel
+ %a_0 = OpVariable %_ptr_Function_v3int Function %9
+ %b_0 = OpVariable %_ptr_Function_int Function %10
+ %20 = OpAccessChain %_ptr_Function_int %a_0 %10
+ %21 = OpLoad %int %20
+ OpStore %b_0 %21
+ %23 = OpAccessChain %_ptr_Function_int %a_0 %10
+ OpStore %23 %int_42
+ OpReturn
+ OpFunctionEnd
+ %deref_let = OpFunction %void None %1
+ %25 = OpLabel
+ %a_1 = OpVariable %_ptr_Function_v3int Function %9
+ %b_1 = OpVariable %_ptr_Function_int Function %10
+ %27 = OpAccessChain %_ptr_Function_int %a_1 %10
+ %28 = OpLoad %int %27
+ OpStore %b_1 %28
+ %30 = OpAccessChain %_ptr_Function_int %a_1 %10
+ OpStore %30 %int_42
+ OpReturn
+ OpFunctionEnd
+%no_deref_let = OpFunction %void None %1
+ %32 = OpLabel
+ %a_2 = OpVariable %_ptr_Function_v3int Function %9
+ %b_2 = OpVariable %_ptr_Function_int Function %10
+ %34 = OpAccessChain %_ptr_Function_int %a_2 %10
+ %35 = OpLoad %int %34
+ OpStore %b_2 %35
+ %37 = OpAccessChain %_ptr_Function_int %a_2 %10
+ OpStore %37 %int_42
+ OpReturn
+ OpFunctionEnd
+ %deref_var = OpFunction %void None %1
+ %39 = OpLabel
+ %a_3 = OpVariable %_ptr_Function_v3int Function %9
+ %i = OpVariable %_ptr_Function_int Function %10
+ %b_3 = OpVariable %_ptr_Function_int Function %10
+ OpStore %i %10
+ %42 = OpLoad %int %i
+ %43 = OpAccessChain %_ptr_Function_int %a_3 %42
+ %44 = OpLoad %int %43
+ OpStore %b_3 %44
+ %46 = OpAccessChain %_ptr_Function_int %a_3 %10
+ OpStore %46 %int_42
+ OpReturn
+ OpFunctionEnd
+%no_deref_var = OpFunction %void None %1
+ %48 = OpLabel
+ %a_4 = OpVariable %_ptr_Function_v3int Function %9
+ %b_4 = OpVariable %_ptr_Function_int Function %10
+ %50 = OpAccessChain %_ptr_Function_int %a_4 %10
+ %51 = OpLoad %int %50
+ OpStore %b_4 %51
+ %53 = OpAccessChain %_ptr_Function_int %a_4 %10
+ OpStore %53 %int_42
+ OpReturn
+ OpFunctionEnd
+ %main = OpFunction %void None %1
+ %55 = OpLabel
+ %56 = OpFunctionCall %void %deref_const
+ %57 = OpFunctionCall %void %no_deref_const
+ %58 = OpFunctionCall %void %deref_let
+ %59 = OpFunctionCall %void %no_deref_let
+ %60 = OpFunctionCall %void %deref_var
+ %61 = OpFunctionCall %void %no_deref_var
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/ptr_sugar/vector_index.wgsl.expected.wgsl b/test/tint/ptr_sugar/vector_index.wgsl.expected.wgsl
new file mode 100644
index 0000000..28f45ae
--- /dev/null
+++ b/test/tint/ptr_sugar/vector_index.wgsl.expected.wgsl
@@ -0,0 +1,55 @@
+fn deref_const() {
+ var a : vec3<i32>;
+ let p = &(a);
+ var b = (*(p))[0];
+ (*(p))[0] = 42;
+}
+
+fn no_deref_const() {
+ var a : vec3<i32>;
+ let p = &(a);
+ var b = p[0];
+ p[0] = 42;
+}
+
+fn deref_let() {
+ var a : vec3<i32>;
+ let p = &(a);
+ let i = 0;
+ var b = (*(p))[i];
+ (*(p))[0] = 42;
+}
+
+fn no_deref_let() {
+ var a : vec3<i32>;
+ let p = &(a);
+ let i = 0;
+ var b = p[i];
+ p[0] = 42;
+}
+
+fn deref_var() {
+ var a : vec3<i32>;
+ let p = &(a);
+ var i = 0;
+ var b = (*(p))[i];
+ (*(p))[0] = 42;
+}
+
+fn no_deref_var() {
+ var a : vec3<i32>;
+ let p = &(a);
+ let i = 0;
+ var b = p[i];
+ p[0] = 42;
+}
+
+@compute @workgroup_size(1)
+fn main() {
+ deref_const();
+ no_deref_const();
+ deref_let();
+ no_deref_let();
+ deref_var();
+ no_deref_var();
+}
diff --git a/test/tint/ptr_sugar/vector_member.wgsl b/test/tint/ptr_sugar/vector_member.wgsl
new file mode 100644
index 0000000..cd2eab8
--- /dev/null
+++ b/test/tint/ptr_sugar/vector_member.wgsl
@@ -0,0 +1,19 @@
+fn deref() {
+ var a : vec3<i32>;
+ let p = &a;
+ var b = (*p).x;
+ (*p).x = 42;
+}
+
+fn no_deref() {
+ var a : vec3<i32>;
+ let p = &a;
+ var b = p.x;
+ p.x = 42;
+}
+
+@compute @workgroup_size(1)
+fn main() {
+ deref();
+ no_deref();
+}
diff --git a/test/tint/ptr_sugar/vector_member.wgsl.expected.dxc.hlsl b/test/tint/ptr_sugar/vector_member.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..b9ff0dc
--- /dev/null
+++ b/test/tint/ptr_sugar/vector_member.wgsl.expected.dxc.hlsl
@@ -0,0 +1,18 @@
+void deref() {
+ int3 a = int3(0, 0, 0);
+ int b = a.x;
+ a.x = 42;
+}
+
+void no_deref() {
+ int3 a = int3(0, 0, 0);
+ int b = a.x;
+ a.x = 42;
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+ deref();
+ no_deref();
+ return;
+}
diff --git a/test/tint/ptr_sugar/vector_member.wgsl.expected.fxc.hlsl b/test/tint/ptr_sugar/vector_member.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..b9ff0dc
--- /dev/null
+++ b/test/tint/ptr_sugar/vector_member.wgsl.expected.fxc.hlsl
@@ -0,0 +1,18 @@
+void deref() {
+ int3 a = int3(0, 0, 0);
+ int b = a.x;
+ a.x = 42;
+}
+
+void no_deref() {
+ int3 a = int3(0, 0, 0);
+ int b = a.x;
+ a.x = 42;
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+ deref();
+ no_deref();
+ return;
+}
diff --git a/test/tint/ptr_sugar/vector_member.wgsl.expected.glsl b/test/tint/ptr_sugar/vector_member.wgsl.expected.glsl
new file mode 100644
index 0000000..4380434
--- /dev/null
+++ b/test/tint/ptr_sugar/vector_member.wgsl.expected.glsl
@@ -0,0 +1,24 @@
+#version 310 es
+
+void deref() {
+ ivec3 a = ivec3(0, 0, 0);
+ int b = a.x;
+ a.x = 42;
+}
+
+void no_deref() {
+ ivec3 a = ivec3(0, 0, 0);
+ int b = a.x;
+ a.x = 42;
+}
+
+void tint_symbol() {
+ deref();
+ no_deref();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ tint_symbol();
+ return;
+}
diff --git a/test/tint/ptr_sugar/vector_member.wgsl.expected.msl b/test/tint/ptr_sugar/vector_member.wgsl.expected.msl
new file mode 100644
index 0000000..b4b88f1
--- /dev/null
+++ b/test/tint/ptr_sugar/vector_member.wgsl.expected.msl
@@ -0,0 +1,21 @@
+#include <metal_stdlib>
+
+using namespace metal;
+void deref() {
+ int3 a = 0;
+ int b = a[0];
+ a[0] = 42;
+}
+
+void no_deref() {
+ int3 a = 0;
+ int b = a[0];
+ a[0] = 42;
+}
+
+kernel void tint_symbol() {
+ deref();
+ no_deref();
+ return;
+}
+
diff --git a/test/tint/ptr_sugar/vector_member.wgsl.expected.spvasm b/test/tint/ptr_sugar/vector_member.wgsl.expected.spvasm
new file mode 100644
index 0000000..90d445f
--- /dev/null
+++ b/test/tint/ptr_sugar/vector_member.wgsl.expected.spvasm
@@ -0,0 +1,55 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 30
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %deref "deref"
+ OpName %a "a"
+ OpName %b "b"
+ OpName %no_deref "no_deref"
+ OpName %a_0 "a"
+ OpName %b_0 "b"
+ OpName %main "main"
+ %void = OpTypeVoid
+ %1 = OpTypeFunction %void
+ %int = OpTypeInt 32 1
+ %v3int = OpTypeVector %int 3
+%_ptr_Function_v3int = OpTypePointer Function %v3int
+ %9 = OpConstantNull %v3int
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+%_ptr_Function_int = OpTypePointer Function %int
+ %16 = OpConstantNull %int
+ %int_42 = OpConstant %int 42
+ %deref = OpFunction %void None %1
+ %4 = OpLabel
+ %a = OpVariable %_ptr_Function_v3int Function %9
+ %b = OpVariable %_ptr_Function_int Function %16
+ %13 = OpAccessChain %_ptr_Function_int %a %uint_0
+ %14 = OpLoad %int %13
+ OpStore %b %14
+ %17 = OpAccessChain %_ptr_Function_int %a %uint_0
+ OpStore %17 %int_42
+ OpReturn
+ OpFunctionEnd
+ %no_deref = OpFunction %void None %1
+ %20 = OpLabel
+ %a_0 = OpVariable %_ptr_Function_v3int Function %9
+ %b_0 = OpVariable %_ptr_Function_int Function %16
+ %22 = OpAccessChain %_ptr_Function_int %a_0 %uint_0
+ %23 = OpLoad %int %22
+ OpStore %b_0 %23
+ %25 = OpAccessChain %_ptr_Function_int %a_0 %uint_0
+ OpStore %25 %int_42
+ OpReturn
+ OpFunctionEnd
+ %main = OpFunction %void None %1
+ %27 = OpLabel
+ %28 = OpFunctionCall %void %deref
+ %29 = OpFunctionCall %void %no_deref
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/ptr_sugar/vector_member.wgsl.expected.wgsl b/test/tint/ptr_sugar/vector_member.wgsl.expected.wgsl
new file mode 100644
index 0000000..4c48f2b
--- /dev/null
+++ b/test/tint/ptr_sugar/vector_member.wgsl.expected.wgsl
@@ -0,0 +1,19 @@
+fn deref() {
+ var a : vec3<i32>;
+ let p = &(a);
+ var b = (*(p)).x;
+ (*(p)).x = 42;
+}
+
+fn no_deref() {
+ var a : vec3<i32>;
+ let p = &(a);
+ var b = p.x;
+ p.x = 42;
+}
+
+@compute @workgroup_size(1)
+fn main() {
+ deref();
+ no_deref();
+}
