1 # Test OpPtrAccessChain applied to workgroup memory.
3 # Derived from the following OpenCL C:
5 #int get_data(local int *d);
7 #int get_data(local int *d) {
11 #kernel void foo(global int *A, global int*B, global int*C) __attribute__((reqd_work_group_size(16,1,1))) {
13 # uint i = get_local_id(0);
14 # data[i] = A[i] * B[i];
16 # if (i == 0) data[16] = 0;
18 # C[i] = get_data(&data[i]);
21 # Compiled with `clspv -no-inline-single -cl-opt-disable <X.clc>` with a (correct) ArrayStride decoration
22 # added. This is allowed but not needed so provides a baseline for comparison.
24 [compute shader spirv]
26 OpCapability VariablePointers
27 OpExtension "SPV_KHR_storage_buffer_storage_class"
28 OpExtension "SPV_KHR_variable_pointers"
29 OpMemoryModel Logical GLSL450
30 OpEntryPoint GLCompute %30 "main" %gl_LocalInvocationID
31 OpExecutionMode %30 LocalSize 16 1 1
33 OpDecorate %_runtimearr_uint ArrayStride 4
34 OpMemberDecorate %_struct_3 0 Offset 0
35 OpDecorate %_struct_3 Block
36 OpDecorate %gl_LocalInvocationID BuiltIn LocalInvocationId
37 OpDecorate %22 DescriptorSet 0
38 OpDecorate %22 Binding 0
39 OpDecorate %23 DescriptorSet 0
40 OpDecorate %23 Binding 1
41 OpDecorate %24 DescriptorSet 0
42 OpDecorate %24 Binding 2
43 OpDecorate %_arr_uint_uint_17 ArrayStride 4
44 OpDecorate %_ptr_Workgroup_uint ArrayStride 4
45 %uint = OpTypeInt 32 0
46 %_runtimearr_uint = OpTypeRuntimeArray %uint
47 %_struct_3 = OpTypeStruct %_runtimearr_uint
48 %_ptr_StorageBuffer__struct_3 = OpTypePointer StorageBuffer %_struct_3
49 %_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
50 %6 = OpTypeFunction %uint %_ptr_Workgroup_uint
52 %8 = OpTypeFunction %void
53 %v3uint = OpTypeVector %uint 3
54 %_ptr_Input_v3uint = OpTypePointer Input %v3uint
55 %_ptr_Input_uint = OpTypePointer Input %uint
56 %_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint
57 %uint_17 = OpConstant %uint 17
58 %_arr_uint_uint_17 = OpTypeArray %uint %uint_17
59 %_ptr_Workgroup__arr_uint_uint_17 = OpTypePointer Workgroup %_arr_uint_uint_17
61 %uint_1 = OpConstant %uint 1
62 %uint_0 = OpConstant %uint 0
63 %uint_16 = OpConstant %uint 16
64 %20 = OpVariable %_ptr_Workgroup__arr_uint_uint_17 Workgroup
65 %gl_LocalInvocationID = OpVariable %_ptr_Input_v3uint Input
66 %22 = OpVariable %_ptr_StorageBuffer__struct_3 StorageBuffer
67 %23 = OpVariable %_ptr_StorageBuffer__struct_3 StorageBuffer
68 %24 = OpVariable %_ptr_StorageBuffer__struct_3 StorageBuffer
69 %25 = OpFunction %uint Pure %6
70 %26 = OpFunctionParameter %_ptr_Workgroup_uint
72 %28 = OpPtrAccessChain %_ptr_Workgroup_uint %26 %uint_1
73 %29 = OpLoad %uint %28
76 %30 = OpFunction %void None %8
78 %32 = OpAccessChain %_ptr_Input_uint %gl_LocalInvocationID %uint_0
79 %33 = OpLoad %uint %32
80 %34 = OpAccessChain %_ptr_StorageBuffer_uint %22 %uint_0 %33
81 %35 = OpLoad %uint %34
82 %36 = OpAccessChain %_ptr_StorageBuffer_uint %23 %uint_0 %33
83 %37 = OpLoad %uint %36
84 %38 = OpIMul %uint %37 %35
85 %39 = OpAccessChain %_ptr_Workgroup_uint %20 %33
87 %40 = OpIEqual %bool %33 %uint_0
88 OpSelectionMerge %43 None
89 OpBranchConditional %40 %41 %43
91 %42 = OpAccessChain %_ptr_Workgroup_uint %20 %uint_16
95 %44 = OpFunctionCall %uint %25 %39
96 %45 = OpAccessChain %_ptr_StorageBuffer_uint %24 %uint_0 %33
103 ssbo 0:0 subdata int 0 0 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15
105 ssbo 0:1 subdata int 0 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1
106 # The answer array C[]
107 ssbo 0:2 subdata int 0 -1 -1 -1 -1 -1 -1 -1 -1 -1 -1 -1 -1 -1 -1 -1 -1
111 probe ssbo int 0:2 0 == 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 0