+# Test OpPtrAccessChain applied to workgroup memory.
+#
+# Derived from the following OpenCL C:
+#
+#int get_data(local int *d);
+#
+#int get_data(local int *d) {
+# return d[1];
+#}
+#
+#kernel void foo(global int *A, global int*B, global int*C) __attribute__((reqd_work_group_size(16,1,1))) {
+# local int data[17];
+# uint i = get_local_id(0);
+# data[i] = A[i] * B[i];
+#
+# if (i == 0) data[16] = 0;
+#
+# C[i] = get_data(&data[i]);
+#}
+#
+# Compiled with `clspv -no-inline-single -cl-opt-disable <X.clc>` with an incorrect ArrayStride decoration
+# added. This decoration should be ignored, so it should give the same results as ArrayStride == 4.
+
+[compute shader spirv]
+ OpCapability Shader
+ OpCapability VariablePointers
+ OpExtension "SPV_KHR_storage_buffer_storage_class"
+ OpExtension "SPV_KHR_variable_pointers"
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %30 "main" %gl_LocalInvocationID
+ OpExecutionMode %30 LocalSize 16 1 1
+ OpSource OpenCL_C 120
+ OpDecorate %_runtimearr_uint ArrayStride 4
+ OpMemberDecorate %_struct_3 0 Offset 0
+ OpDecorate %_struct_3 Block
+ OpDecorate %gl_LocalInvocationID BuiltIn LocalInvocationId
+ OpDecorate %22 DescriptorSet 0
+ OpDecorate %22 Binding 0
+ OpDecorate %23 DescriptorSet 0
+ OpDecorate %23 Binding 1
+ OpDecorate %24 DescriptorSet 0
+ OpDecorate %24 Binding 2
+ OpDecorate %_arr_uint_uint_17 ArrayStride 4
+ OpDecorate %_ptr_Workgroup_uint ArrayStride 8
+ %uint = OpTypeInt 32 0
+%_runtimearr_uint = OpTypeRuntimeArray %uint
+ %_struct_3 = OpTypeStruct %_runtimearr_uint
+%_ptr_StorageBuffer__struct_3 = OpTypePointer StorageBuffer %_struct_3
+%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
+ %6 = OpTypeFunction %uint %_ptr_Workgroup_uint
+ %void = OpTypeVoid
+ %8 = OpTypeFunction %void
+ %v3uint = OpTypeVector %uint 3
+%_ptr_Input_v3uint = OpTypePointer Input %v3uint
+%_ptr_Input_uint = OpTypePointer Input %uint
+%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint
+ %uint_17 = OpConstant %uint 17
+%_arr_uint_uint_17 = OpTypeArray %uint %uint_17
+%_ptr_Workgroup__arr_uint_uint_17 = OpTypePointer Workgroup %_arr_uint_uint_17
+ %bool = OpTypeBool
+ %uint_1 = OpConstant %uint 1
+ %uint_0 = OpConstant %uint 0
+ %uint_16 = OpConstant %uint 16
+ %20 = OpVariable %_ptr_Workgroup__arr_uint_uint_17 Workgroup
+%gl_LocalInvocationID = OpVariable %_ptr_Input_v3uint Input
+ %22 = OpVariable %_ptr_StorageBuffer__struct_3 StorageBuffer
+ %23 = OpVariable %_ptr_StorageBuffer__struct_3 StorageBuffer
+ %24 = OpVariable %_ptr_StorageBuffer__struct_3 StorageBuffer
+ %25 = OpFunction %uint Pure %6
+ %26 = OpFunctionParameter %_ptr_Workgroup_uint
+ %27 = OpLabel
+ %28 = OpPtrAccessChain %_ptr_Workgroup_uint %26 %uint_1
+ %29 = OpLoad %uint %28
+ OpReturnValue %29
+ OpFunctionEnd
+ %30 = OpFunction %void None %8
+ %31 = OpLabel
+ %32 = OpAccessChain %_ptr_Input_uint %gl_LocalInvocationID %uint_0
+ %33 = OpLoad %uint %32
+ %34 = OpAccessChain %_ptr_StorageBuffer_uint %22 %uint_0 %33
+ %35 = OpLoad %uint %34
+ %36 = OpAccessChain %_ptr_StorageBuffer_uint %23 %uint_0 %33
+ %37 = OpLoad %uint %36
+ %38 = OpIMul %uint %37 %35
+ %39 = OpAccessChain %_ptr_Workgroup_uint %20 %33
+ OpStore %39 %38
+ %40 = OpIEqual %bool %33 %uint_0
+ OpSelectionMerge %43 None
+ OpBranchConditional %40 %41 %43
+ %41 = OpLabel
+ %42 = OpAccessChain %_ptr_Workgroup_uint %20 %uint_16
+ OpStore %42 %uint_0
+ OpBranch %43
+ %43 = OpLabel
+ %44 = OpFunctionCall %uint %25 %39
+ %45 = OpAccessChain %_ptr_StorageBuffer_uint %24 %uint_0 %33
+ OpStore %45 %44
+ OpReturn
+ OpFunctionEnd
+
+[test]
+# A[]
+ssbo 0:0 subdata int 0 0 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15
+# B[]
+ssbo 0:1 subdata int 0 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1
+# The answer array C[]
+ssbo 0:2 subdata int 0 -1 -1 -1 -1 -1 -1 -1 -1 -1 -1 -1 -1 -1 -1 -1 -1
+
+compute 1 1 1
+
+probe ssbo int 0:2 0 == 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 0