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
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184
185
186
187
188
189
190
191
192
193
194
195
196
197
198
199
200
201
202
|
// RUN: %clang_cc1 %s -triple spir-unknown-unknown -emit-llvm -pedantic -verify -O0 -o - -DDECL | FileCheck %s --check-prefixes="COMMON,EXPL"
// RUN: %clang_cc1 %s -triple spir-unknown-unknown -emit-llvm -pedantic -verify -O0 -o - -DDECL -DUSE_DEFLT | FileCheck %s --check-prefixes="COMMON,IMPL"
// RUN: %clang_cc1 %s -triple spir-unknown-unknown -emit-llvm -pedantic -verify -O0 -o - | FileCheck %s --check-prefixes="COMMON,IMPL"
// expected-no-diagnostics
// Test that the 'this' pointer is in the __generic address space.
#ifdef USE_DEFLT
#define DEFAULT =default
#else
#define DEFAULT
#endif
class C {
public:
int v;
#ifdef DECL
C() DEFAULT;
C(C &&c) DEFAULT;
C(const C &c) DEFAULT;
C &operator=(const C &c) DEFAULT;
C &operator=(C &&c) & DEFAULT;
#endif
C operator+(const C& c) {
v += c.v;
return *this;
}
int get() { return v; }
int outside();
};
#if defined(DECL) && !defined(USE_DEFLT)
C::C() { v = 2; };
C::C(C &&c) { v = c.v; }
C::C(const C &c) { v = c.v; }
C &C::operator=(const C &c) {
v = c.v;
return *this;
}
C &C::operator=(C &&c) & {
v = c.v;
return *this;
}
#endif
int C::outside() {
return v;
}
extern C&& foo();
__global C c;
__kernel void test__global() {
int i = c.get();
int i2 = (&c)->get();
int i3 = c.outside();
C c1(c);
C c2;
c2 = c1;
C c3 = c1 + c2;
C c4(foo());
C c5 = foo();
}
// Test that the address space is __generic for all members
// EXPL: @_ZNU3AS41CC2Ev(ptr addrspace(4) {{[^,]*}} %this)
// EXPL: @_ZNU3AS41CC1Ev(ptr addrspace(4) {{[^,]*}} %this)
// EXPL: @_ZNU3AS41CC2EOU3AS4S_(ptr addrspace(4) {{[^,]*}} %this
// EXPL: @_ZNU3AS41CC1EOU3AS4S_(ptr addrspace(4) {{[^,]*}} %this
// EXPL: @_ZNU3AS41CC2ERU3AS4KS_(ptr addrspace(4) {{[^,]*}} %this
// EXPL: @_ZNU3AS41CC1ERU3AS4KS_(ptr addrspace(4) {{[^,]*}} %this
// EXPL: @_ZNU3AS41CaSERU3AS4KS_(ptr addrspace(4) {{[^,]*}} %this
// EXPL: @_ZNU3AS4R1CaSEOU3AS4S_(ptr addrspace(4) {{[^,]*}} %this
// COMMON: @_ZNU3AS41C7outsideEv(ptr addrspace(4) {{[^,]*}} %this)
// EXPL-LABEL: @__cxx_global_var_init()
// EXPL: call spir_func void @_ZNU3AS41CC1Ev(ptr addrspace(4) {{[^,]*}} addrspacecast (ptr addrspace(1) @c to ptr addrspace(4)))
// COMMON-LABEL: @test__global()
// Test the address space of 'this' when invoking a method.
// COMMON: call spir_func noundef i32 @_ZNU3AS41C3getEv(ptr addrspace(4) {{[^,]*}} addrspacecast (ptr addrspace(1) @c to ptr addrspace(4)))
// Test the address space of 'this' when invoking a method using a pointer to the object.
// COMMON: call spir_func noundef i32 @_ZNU3AS41C3getEv(ptr addrspace(4) {{[^,]*}} addrspacecast (ptr addrspace(1) @c to ptr addrspace(4)))
// Test the address space of 'this' when invoking a method that is declared in the file contex.
// COMMON: call spir_func noundef i32 @_ZNU3AS41C7outsideEv(ptr addrspace(4) {{[^,]*}} addrspacecast (ptr addrspace(1) @c to ptr addrspace(4)))
// Test the address space of 'this' when invoking copy-constructor.
// COMMON: [[C1GEN:%[.a-z0-9]+]] = addrspacecast ptr %c1 to ptr addrspace(4)
// IMPL: call void @llvm.memcpy.p0.p4.i32(ptr {{.*}}%c1, ptr addrspace(4) {{.*}}addrspacecast (ptr addrspace(1) @c to ptr addrspace(4))
// EXPL: call spir_func void @_ZNU3AS41CC1ERU3AS4KS_(ptr addrspace(4) {{[^,]*}} [[C1GEN]], ptr addrspace(4) noundef align 4 dereferenceable(4) addrspacecast (ptr addrspace(1) @c to ptr addrspace(4)))
// Test the address space of 'this' when invoking a constructor.
// EXPL: [[C2GEN:%[.a-z0-9]+]] = addrspacecast ptr %c2 to ptr addrspace(4)
// EXPL: call spir_func void @_ZNU3AS41CC1Ev(ptr addrspace(4) {{[^,]*}} [[C2GEN]])
// Test the address space of 'this' when invoking assignment operator.
// COMMON: [[C1GEN:%[.a-z0-9]+]] = addrspacecast ptr %c1 to ptr addrspace(4)
// COMMON: [[C2GEN:%[.a-z0-9]+]] = addrspacecast ptr %c2 to ptr addrspace(4)
// EXPL: call spir_func noundef align 4 dereferenceable(4) ptr addrspace(4) @_ZNU3AS41CaSERU3AS4KS_(ptr addrspace(4) {{[^,]*}} [[C2GEN]], ptr addrspace(4) noundef align 4 dereferenceable(4) [[C1GEN]])
// IMPL: call void @llvm.memcpy.p4.p4.i32(ptr addrspace(4) {{.*}}[[C2GEN]], ptr addrspace(4) {{.*}}[[C1GEN]]
// Test the address space of 'this' when invoking the operator+
// COMMON: [[C1GEN:%[.a-z0-9]+]] = addrspacecast ptr %c1 to ptr addrspace(4)
// COMMON: [[C2GEN:%[.a-z0-9]+]] = addrspacecast ptr %c2 to ptr addrspace(4)
// COMMON: call spir_func void @_ZNU3AS41CplERU3AS4KS_(ptr dead_on_unwind writable sret(%class.C) align 4 %c3, ptr addrspace(4) {{[^,]*}} [[C1GEN]], ptr addrspace(4) noundef align 4 dereferenceable(4) [[C2GEN]])
// Test the address space of 'this' when invoking the move constructor
// COMMON: [[C4GEN:%[.a-z0-9]+]] = addrspacecast ptr %c4 to ptr addrspace(4)
// COMMON: [[CALL:%call[0-9]+]] = call spir_func noundef align 4 dereferenceable(4) ptr addrspace(4) @_Z3foov()
// EXPL: call spir_func void @_ZNU3AS41CC1EOU3AS4S_(ptr addrspace(4) {{[^,]*}} [[C4GEN]], ptr addrspace(4) noundef align 4 dereferenceable(4) [[CALL]])
// IMPL: call void @llvm.memcpy.p0.p4.i32(ptr {{.*}}%c4, ptr addrspace(4) {{.*}}[[CALL]]
// Test the address space of 'this' when invoking the move assignment
// COMMON: [[C5GEN:%[.a-z0-9]+]] = addrspacecast ptr %c5 to ptr addrspace(4)
// COMMON: [[CALL:%call[0-9]+]] = call spir_func noundef align 4 dereferenceable(4) ptr addrspace(4) @_Z3foov()
// EXPL: call spir_func void @_ZNU3AS41CC1EOU3AS4S_(ptr addrspace(4) {{[^,]*}} [[C5GEN]], ptr addrspace(4) noundef align 4 dereferenceable(4) [[CALL]])
// IMPL: call void @llvm.memcpy.p0.p4.i32(ptr {{.*}}%c5, ptr addrspace(4) {{.*}}[[CALL]]
// Tests address space of inline members
//COMMON: @_ZNU3AS41C3getEv(ptr addrspace(4) {{[^,]*}} %this)
//COMMON: @_ZNU3AS41CplERU3AS4KS_(ptr dead_on_unwind noalias writable sret(%class.C) align 4 %agg.result, ptr addrspace(4) {{[^,]*}} %this
#define TEST(AS) \
__kernel void test##AS() { \
AS C c; \
int i = c.get(); \
C c1(c); \
C c2; \
c2 = c1; \
}
TEST(__local)
// COMMON-LABEL: @test__local
// Test that we don't initialize an object in local address space.
// EXPL-NOT: call spir_func void @_ZNU3AS41CC1Ev(ptr addrspace(4) addrspacecast (ptr addrspace(3) @_ZZ11test__localE1c to ptr addrspace(4)))
// Test the address space of 'this' when invoking a method.
// COMMON: call spir_func noundef i32 @_ZNU3AS41C3getEv(ptr addrspace(4) {{[^,]*}} addrspacecast (ptr addrspace(3) @_ZZ11test__localE1c to ptr addrspace(4)))
// Test the address space of 'this' when invoking copy-constructor.
// COMMON: [[C1GEN:%[.a-z0-9]+]] = addrspacecast ptr %c1 to ptr addrspace(4)
// EXPL: call spir_func void @_ZNU3AS41CC1ERU3AS4KS_(ptr addrspace(4) {{[^,]*}} [[C1GEN]], ptr addrspace(4) noundef align 4 dereferenceable(4) addrspacecast (ptr addrspace(3) @_ZZ11test__localE1c to ptr addrspace(4)))
// IMPL: call void @llvm.memcpy.p0.p4.i32(ptr {{.*}}%c1, ptr addrspace(4) {{.*}}addrspacecast (ptr addrspace(3) @_ZZ11test__localE1c to ptr addrspace(4)), i32 4, i1 false)
// Test the address space of 'this' when invoking a constructor.
// EXPL: [[C2GEN:%[.a-z0-9]+]] = addrspacecast ptr %c2 to ptr addrspace(4)
// EXPL: call spir_func void @_ZNU3AS41CC1Ev(ptr addrspace(4) {{[^,]*}} [[C2GEN]])
// Test the address space of 'this' when invoking assignment operator.
// COMMON: [[C1GEN:%[.a-z0-9]+]] = addrspacecast ptr %c1 to ptr addrspace(4)
// COMMON: [[C2GEN:%[.a-z0-9]+]] = addrspacecast ptr %c2 to ptr addrspace(4)
// EXPL: call spir_func noundef align 4 dereferenceable(4) ptr addrspace(4) @_ZNU3AS41CaSERU3AS4KS_(ptr addrspace(4) {{[^,]*}} [[C2GEN]], ptr addrspace(4) noundef align 4 dereferenceable(4) [[C1GEN]])
// IMPL: call void @llvm.memcpy.p4.p4.i32(ptr addrspace(4) {{.*}}[[C2GEN]], ptr addrspace(4) {{.*}}[[C1GEN]]
TEST(__private)
// COMMON-LABEL: @test__private
// Test the address space of 'this' when invoking a constructor for an object in non-default address space
// EXPL: [[CGEN:%[.a-z0-9]+]] = addrspacecast ptr %c to ptr addrspace(4)
// EXPL: call spir_func void @_ZNU3AS41CC1Ev(ptr addrspace(4) {{[^,]*}} [[CGEN]])
// Test the address space of 'this' when invoking a method.
// COMMON: [[CGEN:%[.a-z0-9]+]] = addrspacecast ptr %c to ptr addrspace(4)
// COMMON: call spir_func noundef i32 @_ZNU3AS41C3getEv(ptr addrspace(4) {{[^,]*}} [[CGEN]])
// Test the address space of 'this' when invoking a copy-constructor.
// COMMON: [[C1GEN:%[.a-z0-9]+]] = addrspacecast ptr %c1 to ptr addrspace(4)
// COMMON: [[CGEN:%[.a-z0-9]+]] = addrspacecast ptr %c to ptr addrspace(4)
// EXPL: call spir_func void @_ZNU3AS41CC1ERU3AS4KS_(ptr addrspace(4) {{[^,]*}} [[C1GEN]], ptr addrspace(4) noundef align 4 dereferenceable(4) [[CGEN]])
// IMPL: call void @llvm.memcpy.p0.p4.i32(ptr {{.*}}%c1, ptr addrspace(4) {{.*}}[[CGEN]]
// Test the address space of 'this' when invoking a constructor.
// EXPL: [[C2GEN:%[.a-z0-9]+]] = addrspacecast ptr %c2 to ptr addrspace(4)
// EXPL: call spir_func void @_ZNU3AS41CC1Ev(ptr addrspace(4) {{[^,]*}} [[C2GEN]])
// Test the address space of 'this' when invoking a copy-assignment.
// COMMON: [[C1GEN:%[.a-z0-9]+]] = addrspacecast ptr %c1 to ptr addrspace(4)
// COMMON: [[C2GEN:%[.a-z0-9]+]] = addrspacecast ptr %c2 to ptr addrspace(4)
// EXPL: call spir_func noundef align 4 dereferenceable(4) ptr addrspace(4) @_ZNU3AS41CaSERU3AS4KS_(ptr addrspace(4) {{[^,]*}} [[C2GEN]], ptr addrspace(4) noundef align 4 dereferenceable(4) [[C1GEN]])
// IMPL: call void @llvm.memcpy.p4.p4.i32(ptr addrspace(4) {{.*}}[[C2GEN]], ptr addrspace(4) {{.*}}[[C1GEN]]
// Test that calling a const method from a non-const method does not crash Clang.
class ConstAndNonConstMethod {
public:
void DoConst() const {
}
void DoNonConst() {
DoConst();
}
};
|
