Add pointer e2e test
Add an end-to-end test showing various pointer usages in WGSL.
Change-Id: I310e2dcfba5dc9001cce49820e85377ce4340a23
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/164540
Auto-Submit: dan sinclair <dsinclair@chromium.org>
Reviewed-by: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
diff --git a/test/tint/access/ptr.wgsl b/test/tint/access/ptr.wgsl
new file mode 100644
index 0000000..43de544
--- /dev/null
+++ b/test/tint/access/ptr.wgsl
@@ -0,0 +1,56 @@
+@group(0) @binding(0) var<storage, read_write> s: i32;
+
+var<workgroup> g1 : atomic<i32>;
+
+struct S {
+ a: i32,
+ b: i32,
+}
+
+fn accept_ptr_deref_pass_through(val: ptr<function, i32>) -> i32 {
+ return *val + accept_ptr_deref_call_func(val);
+}
+
+fn accept_ptr_to_struct_and_access(val: ptr<function, S>) -> i32 {
+ return (*val).a + (*val).b;
+}
+
+fn accept_ptr_to_struct_access_pass_ptr(val: ptr<function, S>) -> i32 {
+ let b = &((*val).a);
+ *b = 2;
+ return *b;
+}
+
+fn accept_ptr_deref_call_func(val: ptr<function, i32>) -> i32 {
+ return *val + accept_value(*val);
+}
+
+fn accept_value(val: i32) -> i32 {
+ return val;
+}
+
+fn accept_ptr_vec_access_elements(v1: ptr<function, vec3f>) -> i32 {
+ (*v1).x = cross(*v1, *v1).x;
+ return i32((*v1).x);
+}
+
+fn call_builtin_with_mod_scope_ptr() -> i32 {
+ return atomicLoad(&g1);
+}
+
+@compute @workgroup_size(1)
+fn main() {
+ var v1 = 0;
+ var v2 = S();
+ let v3 = &v2;
+ var v4 = vec3f();
+ let t1 = atomicLoad(&g1);
+
+ s = accept_ptr_deref_pass_through(&v1) +
+ accept_ptr_to_struct_and_access(&v2) +
+ accept_ptr_to_struct_and_access(v3) +
+ accept_ptr_vec_access_elements(&v4) +
+ accept_ptr_to_struct_access_pass_ptr(&v2) +
+ call_builtin_with_mod_scope_ptr() +
+ t1;
+}
diff --git a/test/tint/access/ptr.wgsl.expected.dxc.hlsl b/test/tint/access/ptr.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..4391e25
--- /dev/null
+++ b/test/tint/access/ptr.wgsl.expected.dxc.hlsl
@@ -0,0 +1,78 @@
+int tint_ftoi(float v) {
+ return ((v < 2147483520.0f) ? ((v < -2147483648.0f) ? -2147483648 : int(v)) : 2147483647);
+}
+
+RWByteAddressBuffer s : register(u0);
+groupshared int g1;
+
+struct S {
+ int a;
+ int b;
+};
+
+int accept_value(int val) {
+ return val;
+}
+
+int accept_ptr_deref_call_func(inout int val) {
+ const int tint_symbol_2 = val;
+ const int tint_symbol_3 = accept_value(val);
+ return (tint_symbol_2 + tint_symbol_3);
+}
+
+int accept_ptr_deref_pass_through(inout int val) {
+ const int tint_symbol = val;
+ const int tint_symbol_1 = accept_ptr_deref_call_func(val);
+ return (tint_symbol + tint_symbol_1);
+}
+
+int accept_ptr_to_struct_and_access(inout S val) {
+ return (val.a + val.b);
+}
+
+int accept_ptr_to_struct_access_pass_ptr(inout S val) {
+ val.a = 2;
+ return val.a;
+}
+
+int accept_ptr_vec_access_elements(inout float3 v1) {
+ v1.x = cross(v1, v1).x;
+ return tint_ftoi(v1.x);
+}
+
+int call_builtin_with_mod_scope_ptr() {
+ int atomic_result = 0;
+ InterlockedOr(g1, 0, atomic_result);
+ return atomic_result;
+}
+
+struct tint_symbol_11 {
+ uint local_invocation_index : SV_GroupIndex;
+};
+
+void main_inner(uint local_invocation_index) {
+ {
+ int atomic_result_1 = 0;
+ InterlockedExchange(g1, 0, atomic_result_1);
+ }
+ GroupMemoryBarrierWithGroupSync();
+ int v1 = 0;
+ S v2 = (S)0;
+ float3 v4 = (0.0f).xxx;
+ int atomic_result_2 = 0;
+ InterlockedOr(g1, 0, atomic_result_2);
+ const int t1 = atomic_result_2;
+ const int tint_symbol_4 = accept_ptr_deref_pass_through(v1);
+ const int tint_symbol_5 = accept_ptr_to_struct_and_access(v2);
+ const int tint_symbol_6 = accept_ptr_to_struct_and_access(v2);
+ const int tint_symbol_7 = accept_ptr_vec_access_elements(v4);
+ const int tint_symbol_8 = accept_ptr_to_struct_access_pass_ptr(v2);
+ const int tint_symbol_9 = call_builtin_with_mod_scope_ptr();
+ s.Store(0u, asuint(((((((tint_symbol_4 + tint_symbol_5) + tint_symbol_6) + tint_symbol_7) + tint_symbol_8) + tint_symbol_9) + t1)));
+}
+
+[numthreads(1, 1, 1)]
+void main(tint_symbol_11 tint_symbol_10) {
+ main_inner(tint_symbol_10.local_invocation_index);
+ return;
+}
diff --git a/test/tint/access/ptr.wgsl.expected.fxc.hlsl b/test/tint/access/ptr.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..4391e25
--- /dev/null
+++ b/test/tint/access/ptr.wgsl.expected.fxc.hlsl
@@ -0,0 +1,78 @@
+int tint_ftoi(float v) {
+ return ((v < 2147483520.0f) ? ((v < -2147483648.0f) ? -2147483648 : int(v)) : 2147483647);
+}
+
+RWByteAddressBuffer s : register(u0);
+groupshared int g1;
+
+struct S {
+ int a;
+ int b;
+};
+
+int accept_value(int val) {
+ return val;
+}
+
+int accept_ptr_deref_call_func(inout int val) {
+ const int tint_symbol_2 = val;
+ const int tint_symbol_3 = accept_value(val);
+ return (tint_symbol_2 + tint_symbol_3);
+}
+
+int accept_ptr_deref_pass_through(inout int val) {
+ const int tint_symbol = val;
+ const int tint_symbol_1 = accept_ptr_deref_call_func(val);
+ return (tint_symbol + tint_symbol_1);
+}
+
+int accept_ptr_to_struct_and_access(inout S val) {
+ return (val.a + val.b);
+}
+
+int accept_ptr_to_struct_access_pass_ptr(inout S val) {
+ val.a = 2;
+ return val.a;
+}
+
+int accept_ptr_vec_access_elements(inout float3 v1) {
+ v1.x = cross(v1, v1).x;
+ return tint_ftoi(v1.x);
+}
+
+int call_builtin_with_mod_scope_ptr() {
+ int atomic_result = 0;
+ InterlockedOr(g1, 0, atomic_result);
+ return atomic_result;
+}
+
+struct tint_symbol_11 {
+ uint local_invocation_index : SV_GroupIndex;
+};
+
+void main_inner(uint local_invocation_index) {
+ {
+ int atomic_result_1 = 0;
+ InterlockedExchange(g1, 0, atomic_result_1);
+ }
+ GroupMemoryBarrierWithGroupSync();
+ int v1 = 0;
+ S v2 = (S)0;
+ float3 v4 = (0.0f).xxx;
+ int atomic_result_2 = 0;
+ InterlockedOr(g1, 0, atomic_result_2);
+ const int t1 = atomic_result_2;
+ const int tint_symbol_4 = accept_ptr_deref_pass_through(v1);
+ const int tint_symbol_5 = accept_ptr_to_struct_and_access(v2);
+ const int tint_symbol_6 = accept_ptr_to_struct_and_access(v2);
+ const int tint_symbol_7 = accept_ptr_vec_access_elements(v4);
+ const int tint_symbol_8 = accept_ptr_to_struct_access_pass_ptr(v2);
+ const int tint_symbol_9 = call_builtin_with_mod_scope_ptr();
+ s.Store(0u, asuint(((((((tint_symbol_4 + tint_symbol_5) + tint_symbol_6) + tint_symbol_7) + tint_symbol_8) + tint_symbol_9) + t1)));
+}
+
+[numthreads(1, 1, 1)]
+void main(tint_symbol_11 tint_symbol_10) {
+ main_inner(tint_symbol_10.local_invocation_index);
+ return;
+}
diff --git a/test/tint/access/ptr.wgsl.expected.glsl b/test/tint/access/ptr.wgsl.expected.glsl
new file mode 100644
index 0000000..152432a
--- /dev/null
+++ b/test/tint/access/ptr.wgsl.expected.glsl
@@ -0,0 +1,73 @@
+#version 310 es
+
+int tint_ftoi(float v) {
+ return ((v < 2147483520.0f) ? ((v < -2147483648.0f) ? (-2147483647 - 1) : int(v)) : 2147483647);
+}
+
+layout(binding = 0, std430) buffer s_block_ssbo {
+ int inner;
+} s;
+
+shared int g1;
+struct S {
+ int a;
+ int b;
+};
+
+int accept_value(int val) {
+ return val;
+}
+
+int accept_ptr_deref_call_func(inout int val) {
+ int tint_symbol_3 = val;
+ int tint_symbol_4 = accept_value(val);
+ return (tint_symbol_3 + tint_symbol_4);
+}
+
+int accept_ptr_deref_pass_through(inout int val) {
+ int tint_symbol_1 = val;
+ int tint_symbol_2 = accept_ptr_deref_call_func(val);
+ return (tint_symbol_1 + tint_symbol_2);
+}
+
+int accept_ptr_to_struct_and_access(inout S val) {
+ return (val.a + val.b);
+}
+
+int accept_ptr_to_struct_access_pass_ptr(inout S val) {
+ val.a = 2;
+ return val.a;
+}
+
+int accept_ptr_vec_access_elements(inout vec3 v1) {
+ v1.x = cross(v1, v1).x;
+ return tint_ftoi(v1.x);
+}
+
+int call_builtin_with_mod_scope_ptr() {
+ return atomicOr(g1, 0);
+}
+
+void tint_symbol(uint local_invocation_index) {
+ {
+ atomicExchange(g1, 0);
+ }
+ barrier();
+ int v1 = 0;
+ S v2 = S(0, 0);
+ vec3 v4 = vec3(0.0f);
+ int t1 = atomicOr(g1, 0);
+ int tint_symbol_5 = accept_ptr_deref_pass_through(v1);
+ int tint_symbol_6 = accept_ptr_to_struct_and_access(v2);
+ int tint_symbol_7 = accept_ptr_to_struct_and_access(v2);
+ int tint_symbol_8 = accept_ptr_vec_access_elements(v4);
+ int tint_symbol_9 = accept_ptr_to_struct_access_pass_ptr(v2);
+ int tint_symbol_10 = call_builtin_with_mod_scope_ptr();
+ s.inner = ((((((tint_symbol_5 + tint_symbol_6) + tint_symbol_7) + tint_symbol_8) + tint_symbol_9) + tint_symbol_10) + t1);
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ tint_symbol(gl_LocalInvocationIndex);
+ return;
+}
diff --git a/test/tint/access/ptr.wgsl.expected.ir.msl b/test/tint/access/ptr.wgsl.expected.ir.msl
new file mode 100644
index 0000000..2f4757d
--- /dev/null
+++ b/test/tint/access/ptr.wgsl.expected.ir.msl
@@ -0,0 +1,116 @@
+SKIP: FAILED
+
+<dawn>/src/tint/lang/msl/writer/printer/printer.cc:493 internal compiler error: S = struct @align(4) {
+ a:i32 @offset(0)
+ b:i32 @offset(4)
+}
+
+%b1 = block { # root
+ %s:ptr<storage, i32, read_write> = var @binding_point(0, 0)
+ %g1:ptr<workgroup, atomic<i32>, read_write> = var
+}
+
+%d = func(%val:i32):i32 -> %b2 {
+ %b2 = block {
+ ret %val
+ }
+}
+%c = func(%val_1:ptr<function, i32, read_write>):i32 -> %b3 { # %val_1: 'val'
+ %b3 = block {
+ %7:i32 = load %val_1
+ %8:i32 = let %7
+ %9:i32 = load %val_1
+ %10:i32 = call %d, %9
+ %11:i32 = add %8, %10
+ ret %11
+ }
+}
+%a = func(%val_2:ptr<function, i32, read_write>):i32 -> %b4 { # %val_2: 'val'
+ %b4 = block {
+ %14:i32 = load %val_2
+ %15:i32 = let %14
+ %16:i32 = call %c, %val_2
+ %17:i32 = add %15, %16
+ ret %17
+ }
+}
+%z = func():i32 -> %b5 {
+ %b5 = block {
+ %19:i32 = atomicLoad %g1
+ ret %19
+ }
+}
+%y = func(%v1:ptr<function, vec3<f32>, read_write>):i32 -> %b6 {
+ %b6 = block {
+ %22:vec3<f32> = load %v1
+ %23:vec3<f32> = load %v1
+ %24:vec3<f32> = cross %22, %23
+ %25:f32 = access %24, 0u
+ store_vector_element %v1, 0u, %25
+ %26:f32 = load_vector_element %v1, 0u
+ %27:i32 = call %tint_f32_to_i32, %26
+ ret %27
+ }
+}
+%b = func(%val_3:ptr<function, S, read_write>):i32 -> %b7 { # %val_3: 'val'
+ %b7 = block {
+ %31:ptr<function, i32, read_write> = access %val_3, 0u
+ %32:i32 = load %31
+ %33:ptr<function, i32, read_write> = access %val_3, 1u
+ %34:i32 = load %33
+ %35:i32 = add %32, %34
+ ret %35
+ }
+}
+%tint_symbol = @compute @workgroup_size(1, 1, 1) func(%tint_local_index:u32 [@local_invocation_index]):void -> %b8 {
+ %b8 = block {
+ %38:bool = eq %tint_local_index, 0u
+ if %38 [t: %b9] { # if_1
+ %b9 = block { # true
+ %39:void = atomicStore %g1, 0i
+ exit_if # if_1
+ }
+ }
+ %40:void = msl.threadgroup_barrier 4u
+ %v1_1:ptr<function, i32, read_write> = var, 0i # %v1_1: 'v1'
+ %v2:ptr<function, S, read_write> = var, S(0i)
+ %v3:ptr<function, S, read_write> = let %v2
+ %v4:ptr<function, vec3<f32>, read_write> = var, vec3<f32>(0.0f)
+ %45:i32 = atomicLoad %g1
+ %t1:i32 = let %45
+ %47:i32 = call %a, %v1_1
+ %48:i32 = let %47
+ %49:i32 = call %b, %v2
+ %50:i32 = add %48, %49
+ %51:i32 = let %50
+ %52:i32 = call %b, %v3
+ %53:i32 = add %51, %52
+ %54:i32 = let %53
+ %55:i32 = call %z
+ %56:i32 = add %54, %55
+ %57:i32 = add %56, %t1
+ %58:i32 = let %57
+ %59:i32 = call %y, %v4
+ %60:i32 = add %58, %59
+ store %s, %60
+ ret
+ }
+}
+%tint_f32_to_i32 = func(%value:f32):i32 -> %b10 {
+ %b10 = block {
+ %62:i32 = convert %value
+ %63:bool = gte %value, -2147483648.0f
+ %64:i32 = select -2147483648i, %62, %63
+ %65:bool = lte %value, 2147483520.0f
+ %66:i32 = select 2147483647i, %64, %65
+ ret %66
+ }
+}
+
+unhandled variable address space
+********************************************************************
+* The tint shader compiler has encountered an unexpected error. *
+* *
+* Please help us fix this issue by submitting a bug report at *
+* crbug.com/tint with the source program that triggered the bug. *
+********************************************************************
diff --git a/test/tint/access/ptr.wgsl.expected.msl b/test/tint/access/ptr.wgsl.expected.msl
new file mode 100644
index 0000000..3f2f596
--- /dev/null
+++ b/test/tint/access/ptr.wgsl.expected.msl
@@ -0,0 +1,70 @@
+#include <metal_stdlib>
+
+using namespace metal;
+int tint_ftoi(float v) {
+ return select(2147483647, select(int(v), (-2147483647 - 1), (v < -2147483648.0f)), (v < 2147483520.0f));
+}
+
+struct S {
+ int a;
+ int b;
+};
+
+int accept_value(int val) {
+ return val;
+}
+
+int accept_ptr_deref_call_func(thread int* const val) {
+ int const tint_symbol_3 = *(val);
+ int const tint_symbol_4 = accept_value(*(val));
+ return as_type<int>((as_type<uint>(tint_symbol_3) + as_type<uint>(tint_symbol_4)));
+}
+
+int accept_ptr_deref_pass_through(thread int* const val) {
+ int const tint_symbol_1 = *(val);
+ int const tint_symbol_2 = accept_ptr_deref_call_func(val);
+ return as_type<int>((as_type<uint>(tint_symbol_1) + as_type<uint>(tint_symbol_2)));
+}
+
+int accept_ptr_to_struct_and_access(thread S* const val) {
+ return as_type<int>((as_type<uint>((*(val)).a) + as_type<uint>((*(val)).b)));
+}
+
+int accept_ptr_to_struct_access_pass_ptr(thread S* const val) {
+ (*(val)).a = 2;
+ return (*(val)).a;
+}
+
+int accept_ptr_vec_access_elements(thread float3* const v1) {
+ (*(v1))[0] = cross(*(v1), *(v1))[0];
+ return tint_ftoi((*(v1))[0]);
+}
+
+int call_builtin_with_mod_scope_ptr(threadgroup atomic_int* const tint_symbol_11) {
+ return atomic_load_explicit(tint_symbol_11, memory_order_relaxed);
+}
+
+void tint_symbol_inner(uint local_invocation_index, threadgroup atomic_int* const tint_symbol_12, device int* const tint_symbol_13) {
+ {
+ atomic_store_explicit(tint_symbol_12, 0, memory_order_relaxed);
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ int v1 = 0;
+ S v2 = S{};
+ float3 v4 = float3(0.0f);
+ int const t1 = atomic_load_explicit(tint_symbol_12, memory_order_relaxed);
+ int const tint_symbol_5 = accept_ptr_deref_pass_through(&(v1));
+ int const tint_symbol_6 = accept_ptr_to_struct_and_access(&(v2));
+ int const tint_symbol_7 = accept_ptr_to_struct_and_access(&(v2));
+ int const tint_symbol_8 = accept_ptr_vec_access_elements(&(v4));
+ int const tint_symbol_9 = accept_ptr_to_struct_access_pass_ptr(&(v2));
+ int const tint_symbol_10 = call_builtin_with_mod_scope_ptr(tint_symbol_12);
+ *(tint_symbol_13) = as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(tint_symbol_5) + as_type<uint>(tint_symbol_6)))) + as_type<uint>(tint_symbol_7)))) + as_type<uint>(tint_symbol_8)))) + as_type<uint>(tint_symbol_9)))) + as_type<uint>(tint_symbol_10)))) + as_type<uint>(t1)));
+}
+
+kernel void tint_symbol(device int* tint_symbol_15 [[buffer(0)]], uint local_invocation_index [[thread_index_in_threadgroup]]) {
+ threadgroup atomic_int tint_symbol_14;
+ tint_symbol_inner(local_invocation_index, &(tint_symbol_14), tint_symbol_15);
+ return;
+}
+
diff --git a/test/tint/access/ptr.wgsl.expected.spvasm b/test/tint/access/ptr.wgsl.expected.spvasm
new file mode 100644
index 0000000..c8f548f
--- /dev/null
+++ b/test/tint/access/ptr.wgsl.expected.spvasm
@@ -0,0 +1,189 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 138
+; Schema: 0
+ OpCapability Shader
+ %81 = OpExtInstImport "GLSL.std.450"
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main" %local_invocation_index_1
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %local_invocation_index_1 "local_invocation_index_1"
+ OpName %s_block "s_block"
+ OpMemberName %s_block 0 "inner"
+ OpName %s "s"
+ OpName %g1 "g1"
+ OpName %tint_ftoi "tint_ftoi"
+ OpName %v "v"
+ OpName %accept_value "accept_value"
+ OpName %val "val"
+ OpName %accept_ptr_deref_call_func "accept_ptr_deref_call_func"
+ OpName %val_0 "val"
+ OpName %accept_ptr_deref_pass_through "accept_ptr_deref_pass_through"
+ OpName %val_1 "val"
+ OpName %S "S"
+ OpMemberName %S 0 "a"
+ OpMemberName %S 1 "b"
+ OpName %accept_ptr_to_struct_and_access "accept_ptr_to_struct_and_access"
+ OpName %val_2 "val"
+ OpName %accept_ptr_to_struct_access_pass_ptr "accept_ptr_to_struct_access_pass_ptr"
+ OpName %val_3 "val"
+ OpName %accept_ptr_vec_access_elements "accept_ptr_vec_access_elements"
+ OpName %v1 "v1"
+ OpName %call_builtin_with_mod_scope_ptr "call_builtin_with_mod_scope_ptr"
+ OpName %main_inner "main_inner"
+ OpName %local_invocation_index "local_invocation_index"
+ OpName %v1_0 "v1"
+ OpName %v2 "v2"
+ OpName %v4 "v4"
+ OpName %main "main"
+ OpDecorate %local_invocation_index_1 BuiltIn LocalInvocationIndex
+ OpDecorate %s_block Block
+ OpMemberDecorate %s_block 0 Offset 0
+ OpDecorate %s DescriptorSet 0
+ OpDecorate %s Binding 0
+ OpMemberDecorate %S 0 Offset 0
+ OpMemberDecorate %S 1 Offset 4
+ %uint = OpTypeInt 32 0
+%_ptr_Input_uint = OpTypePointer Input %uint
+%local_invocation_index_1 = OpVariable %_ptr_Input_uint Input
+ %int = OpTypeInt 32 1
+ %s_block = OpTypeStruct %int
+%_ptr_StorageBuffer_s_block = OpTypePointer StorageBuffer %s_block
+ %s = OpVariable %_ptr_StorageBuffer_s_block StorageBuffer
+%_ptr_Workgroup_int = OpTypePointer Workgroup %int
+ %g1 = OpVariable %_ptr_Workgroup_int Workgroup
+ %float = OpTypeFloat 32
+ %10 = OpTypeFunction %int %float
+%float_2_14748352e_09 = OpConstant %float 2.14748352e+09
+ %bool = OpTypeBool
+%float_n2_14748365e_09 = OpConstant %float -2.14748365e+09
+%int_n2147483648 = OpConstant %int -2147483648
+%int_2147483647 = OpConstant %int 2147483647
+ %25 = OpTypeFunction %int %int
+%_ptr_Function_int = OpTypePointer Function %int
+ %29 = OpTypeFunction %int %_ptr_Function_int
+ %S = OpTypeStruct %int %int
+%_ptr_Function_S = OpTypePointer Function %S
+ %47 = OpTypeFunction %int %_ptr_Function_S
+ %uint_0 = OpConstant %uint 0
+ %uint_1 = OpConstant %uint 1
+ %int_2 = OpConstant %int 2
+ %v3float = OpTypeVector %float 3
+%_ptr_Function_v3float = OpTypePointer Function %v3float
+ %71 = OpTypeFunction %int %_ptr_Function_v3float
+%_ptr_Function_float = OpTypePointer Function %float
+ %91 = OpTypeFunction %int
+ %uint_2 = OpConstant %uint 2
+ %void = OpTypeVoid
+ %97 = OpTypeFunction %void %uint
+ %104 = OpConstantNull %int
+ %uint_264 = OpConstant %uint 264
+ %108 = OpConstantNull %S
+ %110 = OpConstantNull %v3float
+%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int
+ %133 = OpTypeFunction %void
+ %tint_ftoi = OpFunction %int None %10
+ %v = OpFunctionParameter %float
+ %14 = OpLabel
+ %17 = OpFOrdLessThan %bool %v %float_2_14748352e_09
+ %21 = OpFOrdLessThan %bool %v %float_n2_14748365e_09
+ %23 = OpConvertFToS %int %v
+ %19 = OpSelect %int %21 %int_n2147483648 %23
+ %15 = OpSelect %int %17 %19 %int_2147483647
+ OpReturnValue %15
+ OpFunctionEnd
+%accept_value = OpFunction %int None %25
+ %val = OpFunctionParameter %int
+ %28 = OpLabel
+ OpReturnValue %val
+ OpFunctionEnd
+%accept_ptr_deref_call_func = OpFunction %int None %29
+ %val_0 = OpFunctionParameter %_ptr_Function_int
+ %33 = OpLabel
+ %35 = OpLoad %int %val_0
+ %38 = OpLoad %int %val_0
+ %36 = OpFunctionCall %int %accept_value %38
+ %39 = OpIAdd %int %35 %36
+ OpReturnValue %39
+ OpFunctionEnd
+%accept_ptr_deref_pass_through = OpFunction %int None %29
+ %val_1 = OpFunctionParameter %_ptr_Function_int
+ %42 = OpLabel
+ %44 = OpLoad %int %val_1
+ %45 = OpFunctionCall %int %accept_ptr_deref_call_func %val_1
+ %46 = OpIAdd %int %44 %45
+ OpReturnValue %46
+ OpFunctionEnd
+%accept_ptr_to_struct_and_access = OpFunction %int None %47
+ %val_2 = OpFunctionParameter %_ptr_Function_S
+ %52 = OpLabel
+ %55 = OpAccessChain %_ptr_Function_int %val_2 %uint_0
+ %56 = OpLoad %int %55
+ %59 = OpAccessChain %_ptr_Function_int %val_2 %uint_1
+ %60 = OpLoad %int %59
+ %61 = OpIAdd %int %56 %60
+ OpReturnValue %61
+ OpFunctionEnd
+%accept_ptr_to_struct_access_pass_ptr = OpFunction %int None %47
+ %val_3 = OpFunctionParameter %_ptr_Function_S
+ %64 = OpLabel
+ %66 = OpAccessChain %_ptr_Function_int %val_3 %uint_0
+ OpStore %66 %int_2
+ %69 = OpAccessChain %_ptr_Function_int %val_3 %uint_0
+ %70 = OpLoad %int %69
+ OpReturnValue %70
+ OpFunctionEnd
+%accept_ptr_vec_access_elements = OpFunction %int None %71
+ %v1 = OpFunctionParameter %_ptr_Function_v3float
+ %76 = OpLabel
+ %79 = OpAccessChain %_ptr_Function_float %v1 %uint_0
+ %83 = OpLoad %v3float %v1
+ %85 = OpLoad %v3float %v1
+ %80 = OpExtInst %v3float %81 Cross %83 %85
+ %86 = OpCompositeExtract %float %80 0
+ OpStore %79 %86
+ %89 = OpAccessChain %_ptr_Function_float %v1 %uint_0
+ %90 = OpLoad %float %89
+ %87 = OpFunctionCall %int %tint_ftoi %90
+ OpReturnValue %87
+ OpFunctionEnd
+%call_builtin_with_mod_scope_ptr = OpFunction %int None %91
+ %93 = OpLabel
+ %94 = OpAtomicLoad %int %g1 %uint_2 %uint_0
+ OpReturnValue %94
+ OpFunctionEnd
+ %main_inner = OpFunction %void None %97
+%local_invocation_index = OpFunctionParameter %uint
+ %101 = OpLabel
+ %v1_0 = OpVariable %_ptr_Function_int Function %104
+ %v2 = OpVariable %_ptr_Function_S Function %108
+ %v4 = OpVariable %_ptr_Function_v3float Function %110
+ OpAtomicStore %g1 %uint_2 %uint_0 %104
+ OpControlBarrier %uint_2 %uint_2 %uint_264
+ OpStore %v1_0 %104
+ OpStore %v2 %108
+ OpStore %v4 %110
+ %112 = OpAtomicLoad %int %g1 %uint_2 %uint_0
+ %114 = OpFunctionCall %int %accept_ptr_deref_pass_through %v1_0
+ %116 = OpFunctionCall %int %accept_ptr_to_struct_and_access %v2
+ %118 = OpFunctionCall %int %accept_ptr_to_struct_and_access %v2
+ %120 = OpFunctionCall %int %accept_ptr_vec_access_elements %v4
+ %122 = OpFunctionCall %int %accept_ptr_to_struct_access_pass_ptr %v2
+ %124 = OpFunctionCall %int %call_builtin_with_mod_scope_ptr
+ %126 = OpAccessChain %_ptr_StorageBuffer_int %s %uint_0
+ %127 = OpIAdd %int %114 %116
+ %128 = OpIAdd %int %127 %118
+ %129 = OpIAdd %int %128 %120
+ %130 = OpIAdd %int %129 %122
+ %131 = OpIAdd %int %130 %124
+ %132 = OpIAdd %int %131 %112
+ OpStore %126 %132
+ OpReturn
+ OpFunctionEnd
+ %main = OpFunction %void None %133
+ %135 = OpLabel
+ %137 = OpLoad %uint %local_invocation_index_1
+ %136 = OpFunctionCall %void %main_inner %137
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/access/ptr.wgsl.expected.wgsl b/test/tint/access/ptr.wgsl.expected.wgsl
new file mode 100644
index 0000000..00fe4bc
--- /dev/null
+++ b/test/tint/access/ptr.wgsl.expected.wgsl
@@ -0,0 +1,49 @@
+@group(0) @binding(0) var<storage, read_write> s : i32;
+
+var<workgroup> g1 : atomic<i32>;
+
+struct S {
+ a : i32,
+ b : i32,
+}
+
+fn accept_ptr_deref_pass_through(val : ptr<function, i32>) -> i32 {
+ return (*(val) + accept_ptr_deref_call_func(val));
+}
+
+fn accept_ptr_to_struct_and_access(val : ptr<function, S>) -> i32 {
+ return ((*(val)).a + (*(val)).b);
+}
+
+fn accept_ptr_to_struct_access_pass_ptr(val : ptr<function, S>) -> i32 {
+ let b = &((*(val)).a);
+ *(b) = 2;
+ return *(b);
+}
+
+fn accept_ptr_deref_call_func(val : ptr<function, i32>) -> i32 {
+ return (*(val) + accept_value(*(val)));
+}
+
+fn accept_value(val : i32) -> i32 {
+ return val;
+}
+
+fn accept_ptr_vec_access_elements(v1 : ptr<function, vec3f>) -> i32 {
+ (*(v1)).x = cross(*(v1), *(v1)).x;
+ return i32((*(v1)).x);
+}
+
+fn call_builtin_with_mod_scope_ptr() -> i32 {
+ return atomicLoad(&(g1));
+}
+
+@compute @workgroup_size(1)
+fn main() {
+ var v1 = 0;
+ var v2 = S();
+ let v3 = &(v2);
+ var v4 = vec3f();
+ let t1 = atomicLoad(&(g1));
+ s = ((((((accept_ptr_deref_pass_through(&(v1)) + accept_ptr_to_struct_and_access(&(v2))) + accept_ptr_to_struct_and_access(v3)) + accept_ptr_vec_access_elements(&(v4))) + accept_ptr_to_struct_access_pass_ptr(&(v2))) + call_builtin_with_mod_scope_ptr()) + t1);
+}
