1// RUN: %clang_cc1 %s -triple spir-unknown-unknown -cl-std=clc++ -emit-llvm -pedantic -verify -O0 -o - -DDECL | FileCheck %s --check-prefixes="COMMON,EXPL"
2// RUN: %clang_cc1 %s -triple spir-unknown-unknown -cl-std=clc++ -emit-llvm -pedantic -verify -O0 -o - -DDECL -DUSE_DEFLT | FileCheck %s --check-prefixes="COMMON,IMPL"
3// RUN: %clang_cc1 %s -triple spir-unknown-unknown -cl-std=clc++ -emit-llvm -pedantic -verify -O0 -o - | FileCheck %s --check-prefixes="COMMON,IMPL"
4// expected-no-diagnostics
5
6// Test that the 'this' pointer is in the __generic address space.
7
8#ifdef USE_DEFLT
9#define DEFAULT =default
10#else
11#define DEFAULT
12#endif
13
14class C {
15public:
16  int v;
17#ifdef DECL
18  C() DEFAULT;
19  C(C &&c) DEFAULT;
20  C(const C &c) DEFAULT;
21  C &operator=(const C &c) DEFAULT;
22  C &operator=(C &&c) & DEFAULT;
23#endif
24  C operator+(const C& c) {
25    v += c.v;
26    return *this;
27  }
28
29  int get() { return v; }
30
31  int outside();
32};
33
34#if defined(DECL) && !defined(USE_DEFLT)
35C::C() { v = 2; };
36
37C::C(C &&c) { v = c.v; }
38
39C::C(const C &c) { v = c.v; }
40
41C &C::operator=(const C &c) {
42  v = c.v;
43  return *this;
44}
45
46C &C::operator=(C &&c) & {
47  v = c.v;
48  return *this;
49}
50#endif
51
52int C::outside() {
53  return v;
54}
55
56extern C&& foo();
57
58__global C c;
59
60__kernel void test__global() {
61  int i = c.get();
62  int i2 = (&c)->get();
63  int i3 = c.outside();
64  C c1(c);
65  C c2;
66  c2 = c1;
67  C c3 = c1 + c2;
68  C c4(foo());
69  C c5 = foo();
70}
71
72// Test that the address space is __generic for all members
73// EXPL: @_ZNU3AS41CC2Ev(%class.C addrspace(4)* {{[^,]*}} %this)
74// EXPL: @_ZNU3AS41CC1Ev(%class.C addrspace(4)* {{[^,]*}} %this)
75// EXPL: @_ZNU3AS41CC2EOU3AS4S_(%class.C addrspace(4)* {{[^,]*}} %this
76// EXPL: @_ZNU3AS41CC1EOU3AS4S_(%class.C addrspace(4)* {{[^,]*}} %this
77// EXPL: @_ZNU3AS41CC2ERU3AS4KS_(%class.C addrspace(4)* {{[^,]*}} %this
78// EXPL: @_ZNU3AS41CC1ERU3AS4KS_(%class.C addrspace(4)* {{[^,]*}} %this
79// EXPL: @_ZNU3AS41CaSERU3AS4KS_(%class.C addrspace(4)* {{[^,]*}} %this
80// EXPL: @_ZNU3AS4R1CaSEOU3AS4S_(%class.C addrspace(4)* {{[^,]*}} %this
81// COMMON: @_ZNU3AS41C7outsideEv(%class.C addrspace(4)* {{[^,]*}} %this)
82
83// EXPL-LABEL: @__cxx_global_var_init()
84// EXPL: call spir_func void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* {{[^,]*}} addrspacecast (%class.C addrspace(1)* @c to %class.C addrspace(4)*))
85
86// COMMON-LABEL: @test__global()
87
88// Test the address space of 'this' when invoking a method.
89// COMMON: call spir_func i32 @_ZNU3AS41C3getEv(%class.C addrspace(4)* {{[^,]*}} addrspacecast (%class.C addrspace(1)* @c to %class.C addrspace(4)*))
90// Test the address space of 'this' when invoking a method using a pointer to the object.
91// COMMON: call spir_func i32 @_ZNU3AS41C3getEv(%class.C addrspace(4)* {{[^,]*}} addrspacecast (%class.C addrspace(1)* @c to %class.C addrspace(4)*))
92
93// Test the address space of 'this' when invoking a method that is declared in the file contex.
94// COMMON: call spir_func i32 @_ZNU3AS41C7outsideEv(%class.C addrspace(4)* {{[^,]*}} addrspacecast (%class.C addrspace(1)* @c to %class.C addrspace(4)*))
95
96// Test the address space of 'this' when invoking copy-constructor.
97// COMMON: [[C1GEN:%[.a-z0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
98// IMPL: [[C1VOID:%[0-9]+]] = bitcast %class.C* %c1 to i8*
99// IMPL: call void @llvm.memcpy.p0i8.p4i8.i32(i8* {{.*}}[[C1VOID]], i8 addrspace(4)* {{.*}}addrspacecast (i8 addrspace(1)* bitcast (%class.C addrspace(1)* @c to i8 addrspace(1)*) to i8 addrspace(4)*)
100// EXPL: call spir_func void @_ZNU3AS41CC1ERU3AS4KS_(%class.C addrspace(4)* {{[^,]*}} [[C1GEN]], %class.C addrspace(4)* align 4 dereferenceable(4) addrspacecast (%class.C addrspace(1)* @c to %class.C addrspace(4)*))
101
102// Test the address space of 'this' when invoking a constructor.
103// EXPL:   [[C2GEN:%[.a-z0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
104// EXPL:   call spir_func void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* {{[^,]*}} [[C2GEN]])
105
106// Test the address space of 'this' when invoking assignment operator.
107// COMMON:  [[C1GEN:%[.a-z0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
108// COMMON:  [[C2GEN:%[.a-z0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
109// EXPL: call spir_func align 4 dereferenceable(4) %class.C addrspace(4)* @_ZNU3AS41CaSERU3AS4KS_(%class.C addrspace(4)* {{[^,]*}} [[C2GEN]], %class.C addrspace(4)* align 4 dereferenceable(4) [[C1GEN]])
110// IMPL: [[C2GENVOID:%[0-9]+]] = bitcast %class.C addrspace(4)* [[C2GEN]] to i8 addrspace(4)*
111// IMPL: [[C1GENVOID:%[0-9]+]] = bitcast %class.C addrspace(4)* [[C1GEN]] to i8 addrspace(4)*
112// IMPL: call void @llvm.memcpy.p4i8.p4i8.i32(i8 addrspace(4)* {{.*}}[[C2GENVOID]], i8 addrspace(4)* {{.*}}[[C1GENVOID]]
113
114// Test the address space of 'this' when invoking the operator+
115// COMMON: [[C1GEN:%[.a-z0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
116// COMMON: [[C2GEN:%[.a-z0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
117// COMMON: call spir_func void @_ZNU3AS41CplERU3AS4KS_(%class.C* sret(%class.C) align 4 %c3, %class.C addrspace(4)* {{[^,]*}} [[C1GEN]], %class.C addrspace(4)* align 4 dereferenceable(4) [[C2GEN]])
118
119// Test the address space of 'this' when invoking the move constructor
120// COMMON: [[C4GEN:%[.a-z0-9]+]] = addrspacecast %class.C* %c4 to %class.C addrspace(4)*
121// COMMON: [[CALL:%call[0-9]+]] = call spir_func align 4 dereferenceable(4) %class.C addrspace(4)* @_Z3foov()
122// EXPL: call spir_func void @_ZNU3AS41CC1EOU3AS4S_(%class.C addrspace(4)* {{[^,]*}} [[C4GEN]], %class.C addrspace(4)* align 4 dereferenceable(4) [[CALL]])
123// IMPL: [[C4VOID:%[0-9]+]] = bitcast %class.C* %c4 to i8*
124// IMPL: [[CALLVOID:%[0-9]+]] = bitcast %class.C addrspace(4)* [[CALL]] to i8 addrspace(4)*
125// IMPL:  call void @llvm.memcpy.p0i8.p4i8.i32(i8* {{.*}}[[C4VOID]], i8 addrspace(4)* {{.*}}[[CALLVOID]]
126
127// Test the address space of 'this' when invoking the move assignment
128// COMMON: [[C5GEN:%[.a-z0-9]+]] = addrspacecast %class.C* %c5 to %class.C addrspace(4)*
129// COMMON: [[CALL:%call[0-9]+]] = call spir_func align 4 dereferenceable(4) %class.C addrspace(4)* @_Z3foov()
130// EXPL: call spir_func void @_ZNU3AS41CC1EOU3AS4S_(%class.C addrspace(4)* {{[^,]*}} [[C5GEN]], %class.C addrspace(4)* align 4 dereferenceable(4) [[CALL]])
131// IMPL: [[C5VOID:%[0-9]+]] = bitcast %class.C* %c5 to i8*
132// IMPL: [[CALLVOID:%[0-9]+]] = bitcast %class.C addrspace(4)* [[CALL]] to i8 addrspace(4)*
133// IMPL: call void @llvm.memcpy.p0i8.p4i8.i32(i8* {{.*}}[[C5VOID]], i8 addrspace(4)* {{.*}}[[CALLVOID]]
134
135// Tests address space of inline members
136//COMMON: @_ZNU3AS41C3getEv(%class.C addrspace(4)* {{[^,]*}} %this)
137//COMMON: @_ZNU3AS41CplERU3AS4KS_(%class.C* noalias sret(%class.C) align 4 %agg.result, %class.C addrspace(4)* {{[^,]*}} %this
138#define TEST(AS)             \
139  __kernel void test##AS() { \
140    AS C c;                  \
141    int i = c.get();         \
142    C c1(c);                 \
143    C c2;                    \
144    c2 = c1;                 \
145  }
146
147TEST(__local)
148
149// COMMON-LABEL: @test__local
150
151// Test that we don't initialize an object in local address space.
152// EXPL-NOT: call spir_func void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* addrspacecast (%class.C addrspace(3)* @_ZZ11test__localE1c to %class.C addrspace(4)*))
153
154// Test the address space of 'this' when invoking a method.
155// COMMON: call spir_func i32 @_ZNU3AS41C3getEv(%class.C addrspace(4)* {{[^,]*}} addrspacecast (%class.C addrspace(3)* @_ZZ11test__localE1c to %class.C addrspace(4)*))
156
157// Test the address space of 'this' when invoking copy-constructor.
158// COMMON: [[C1GEN:%[.a-z0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
159// EXPL: call spir_func void @_ZNU3AS41CC1ERU3AS4KS_(%class.C addrspace(4)* {{[^,]*}} [[C1GEN]], %class.C addrspace(4)* align 4 dereferenceable(4) addrspacecast (%class.C addrspace(3)* @_ZZ11test__localE1c to %class.C addrspace(4)*))
160// IMPL:  [[C1VOID:%[0-9]+]] = bitcast %class.C* %c1 to i8*
161// IMPL:  call void @llvm.memcpy.p0i8.p4i8.i32(i8* {{.*}}[[C1VOID]], i8 addrspace(4)* {{.*}}addrspacecast (i8 addrspace(3)* bitcast (%class.C addrspace(3)* @_ZZ11test__localE1c to i8 addrspace(3)*) to i8 addrspace(4)*), i32 4, i1 false)
162
163// Test the address space of 'this' when invoking a constructor.
164// EXPL: [[C2GEN:%[.a-z0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
165// EXPL: call spir_func void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* {{[^,]*}} [[C2GEN]])
166
167// Test the address space of 'this' when invoking assignment operator.
168// COMMON: [[C1GEN:%[.a-z0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
169// COMMON: [[C2GEN:%[.a-z0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
170// EXPL: call spir_func align 4 dereferenceable(4) %class.C addrspace(4)* @_ZNU3AS41CaSERU3AS4KS_(%class.C addrspace(4)* {{[^,]*}} [[C2GEN]], %class.C addrspace(4)* align 4 dereferenceable(4) [[C1GEN]])
171// IMPL: [[C2GENVOID:%[0-9]+]] = bitcast %class.C addrspace(4)* [[C2GEN]] to i8 addrspace(4)*
172// IMPL: [[C1GENVOID:%[0-9]+]] = bitcast %class.C addrspace(4)* [[C1GEN]] to i8 addrspace(4)*
173// IMPL: call void @llvm.memcpy.p4i8.p4i8.i32(i8 addrspace(4)* {{.*}}[[C2GENVOID]], i8 addrspace(4)* {{.*}}[[C1GENVOID]]
174
175TEST(__private)
176
177// COMMON-LABEL: @test__private
178
179// Test the address space of 'this' when invoking a constructor for an object in non-default address space
180// EXPL: [[CGEN:%[.a-z0-9]+]] = addrspacecast %class.C* %c to %class.C addrspace(4)*
181// EXPL: call spir_func void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* {{[^,]*}} [[CGEN]])
182
183// Test the address space of 'this' when invoking a method.
184// COMMON: [[CGEN:%[.a-z0-9]+]] = addrspacecast %class.C* %c to %class.C addrspace(4)*
185// COMMON: call spir_func i32 @_ZNU3AS41C3getEv(%class.C addrspace(4)* {{[^,]*}} [[CGEN]])
186
187// Test the address space of 'this' when invoking a copy-constructor.
188// COMMON: [[C1GEN:%[.a-z0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
189// COMMON: [[CGEN:%[.a-z0-9]+]] = addrspacecast %class.C* %c to %class.C addrspace(4)*
190// EXPL: call spir_func void @_ZNU3AS41CC1ERU3AS4KS_(%class.C addrspace(4)* {{[^,]*}} [[C1GEN]], %class.C addrspace(4)* align 4 dereferenceable(4) [[CGEN]])
191// IMPL: [[C1VOID:%[0-9]+]] = bitcast %class.C* %c1 to i8*
192// IMPL: [[CGENVOID:%[0-9]+]] = bitcast %class.C addrspace(4)* [[CGEN]] to i8 addrspace(4)*
193// IMPL: call void @llvm.memcpy.p0i8.p4i8.i32(i8* {{.*}}[[C1VOID]], i8 addrspace(4)* {{.*}}[[CGENVOID]]
194
195// Test the address space of 'this' when invoking a constructor.
196// EXPL: [[C2GEN:%[.a-z0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
197// EXPL: call spir_func void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* {{[^,]*}} [[C2GEN]])
198
199// Test the address space of 'this' when invoking a copy-assignment.
200// COMMON: [[C1GEN:%[.a-z0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
201// COMMON: [[C2GEN:%[.a-z0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
202// EXPL: call spir_func align 4 dereferenceable(4) %class.C addrspace(4)* @_ZNU3AS41CaSERU3AS4KS_(%class.C addrspace(4)* {{[^,]*}} [[C2GEN]], %class.C addrspace(4)* align 4 dereferenceable(4) [[C1GEN]])
203// IMPL: [[C2GENVOID:%[0-9]+]] = bitcast %class.C addrspace(4)* [[C2GEN]] to i8 addrspace(4)*
204// IMPL: [[C1GENVOID:%[0-9]+]] = bitcast %class.C addrspace(4)* [[C1GEN]] to i8 addrspace(4)*
205// IMPL:  call void @llvm.memcpy.p4i8.p4i8.i32(i8 addrspace(4)* {{.*}}[[C2GENVOID]], i8 addrspace(4)* {{.*}}[[C1GENVOID]]
206
207// Test that calling a const method from a non-const method does not crash Clang.
208class ConstAndNonConstMethod {
209public:
210  void DoConst() const {
211  }
212
213  void DoNonConst() {
214    DoConst();
215  }
216};
217