1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
|
; Generated from:
;
; kernel void test(read_only image2d_t image, global int* out)
; {
; size_t gid = get_global_id(0);
; const sampler_t sampler = CLK_NORMALIZED_COORDS_TRUE |
; CLK_ADDRESS_REPEAT |
; CLK_FILTER_NEAREST;
;
; if(gid % 2)
; {
; float4 result = read_imagef( image, sampler, (int2)(0, 0));
; out[0] = result.x;
; }
; else
; {
; int4 result = read_imagei( image, sampler, (int2)(1, 1));
; out[0] = result.x;
; }
; }
119734787 65536 393230 42 0
2 Capability Addresses
2 Capability Linkage
2 Capability Kernel
2 Capability ImageBasic
2 Capability LiteralSampler
5 ExtInstImport 1 "OpenCL.std"
3 MemoryModel 1 2
5 EntryPoint 6 10 "test"
3 Source 3 200000
13 Decorate 5 LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import
3 Decorate 5 Constant
4 Decorate 5 BuiltIn 28
4 Decorate 12 FuncParamAttr 5
4 TypeInt 2 32 0
4 Constant 2 21 1
4 Constant 2 23 0
4 TypeVector 3 2 3
4 TypePointer 4 1 3
2 TypeVoid 6
10 TypeImage 7 6 1 0 0 0 0 0 0
4 TypePointer 8 5 2
5 TypeFunction 9 6 7 8
2 TypeSampler 19
2 TypeBool 24
3 TypeSampledImage 26 7
3 TypeFloat 28 32
4 TypeVector 29 28 4
4 TypeVector 30 2 2
4 TypeVector 37 2 4
4 Variable 4 5 1
6 ConstantSampler 19 20 3 1 0
3 ConstantNull 30 31
4 Constant 28 32 0
5 ConstantComposite 30 38 21 21
5 Function 6 10 0 9
3 FunctionParameter 7 11
3 FunctionParameter 8 12
2 Label 13
6 Load 3 17 5 2 16
5 CompositeExtract 2 18 17 0
5 BitwiseAnd 2 22 18 21
5 IEqual 24 25 22 23
4 BranchConditional 25 15 14
2 Label 14
5 SampledImage 26 27 11 20
7 ImageSampleExplicitLod 29 33 27 31 2 32
5 CompositeExtract 28 34 33 0
4 ConvertFToS 2 35 34
2 Branch 16
2 Label 15
5 SampledImage 26 36 11 20
7 ImageSampleExplicitLod 37 39 36 38 2 32
5 CompositeExtract 2 40 39 0
2 Branch 16
2 Label 16
7 Phi 2 41 40 15 35 14
5 Store 12 41 2 4
1 Return
1 FunctionEnd
; RUN: llvm-spirv %s -to-binary -o %t.spv
; RUN: llvm-spirv -r --spirv-target-env=SPV-IR %t.spv -o %t.bc
; RUN: llvm-dis < %t.bc | FileCheck %s --check-prefix=CHECK-LLVM
; Check whether seperate initilizer has been generated for each use
; CHECK-LLVM: %[[cond:.*]] = icmp eq i32 %10, 0
; CHECK-LLVM: br i1 %[[cond]], label %[[br1:.*]], label %[[br0:.*]]
; CHECK-LLVM: [[br0]]:
; CHECK-LLVM: %[[s0:.*]] = call %spirv.Sampler addrspace(2)* @__translate_sampler_initializer(i32 23)
; CHECK-LLVM: call spir_func %spirv.SampledImage._void_1_0_0_0_0_0_0 addrspace(1)* @_Z20__spirv_SampledImagePU3AS133__spirv_Image__void_1_0_0_0_0_0_0PU3AS215__spirv_Sampler(%spirv.Image._void_1_0_0_0_0_0_0 addrspace(1)* %{{.*}}, %spirv.Sampler addrspace(2)* %[[s0]])
; CHECK-LLVM: [[br1]]:
; CHECK-LLVM: %[[s1:.*]] = call %spirv.Sampler addrspace(2)* @__translate_sampler_initializer(i32 23)
; CHECK-LLVM: call spir_func %spirv.SampledImage._void_1_0_0_0_0_0_0 addrspace(1)* @_Z20__spirv_SampledImagePU3AS133__spirv_Image__void_1_0_0_0_0_0_0PU3AS215__spirv_Sampler(%spirv.Image._void_1_0_0_0_0_0_0 addrspace(1)* %{{.*}}, %spirv.Sampler addrspace(2)* %[[s1]])
|
