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);
+}