1 // RUN: %clang_cc1 -no-opaque-pointers %s -triple spir-unknown-unknown -emit-llvm -pedantic -verify -O0 -o - -DDECL | FileCheck %s --check-prefixes="COMMON,EXPL"
2 // RUN: %clang_cc1 -no-opaque-pointers %s -triple spir-unknown-unknown -emit-llvm -pedantic -verify -O0 -o - -DDECL -DUSE_DEFLT | FileCheck %s --check-prefixes="COMMON,IMPL"
3 // RUN: %clang_cc1 -no-opaque-pointers %s -triple spir-unknown-unknown -emit-llvm -pedantic -verify -O0 -o - | FileCheck %s --check-prefixes="COMMON,IMPL"
4 // expected-no-diagnostics
6 // Test that the 'this' pointer is in the __generic address space.
9 #define DEFAULT =default
20 C(const C &c) DEFAULT;
21 C &operator=(const C &c) DEFAULT;
22 C &operator=(C &&c) & DEFAULT;
24 C operator+(const C& c) {
29 int get() { return v; }
34 #if defined(DECL) && !defined(USE_DEFLT)
37 C::C(C &&c) { v = c.v; }
39 C::C(const C &c) { v = c.v; }
41 C &C::operator=(const C &c) {
46 C &C::operator=(C &&c) & {
60 __kernel void test__global() {
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)
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)*))
86 // COMMON-LABEL: @test__global()
88 // Test the address space of 'this' when invoking a method.
89 // COMMON: call spir_func noundef 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 noundef i32 @_ZNU3AS41C3getEv(%class.C addrspace(4)* {{[^,]*}} addrspacecast (%class.C addrspace(1)* @c to %class.C addrspace(4)*))
93 // Test the address space of 'this' when invoking a method that is declared in the file contex.
94 // COMMON: call spir_func noundef i32 @_ZNU3AS41C7outsideEv(%class.C addrspace(4)* {{[^,]*}} addrspacecast (%class.C addrspace(1)* @c to %class.C addrspace(4)*))
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)* noundef align 4 dereferenceable(4) addrspacecast (%class.C addrspace(1)* @c to %class.C addrspace(4)*))
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]])
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 noundef align 4 dereferenceable(4) %class.C addrspace(4)* @_ZNU3AS41CaSERU3AS4KS_(%class.C addrspace(4)* {{[^,]*}} [[C2GEN]], %class.C addrspace(4)* noundef 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]]
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)* noundef align 4 dereferenceable(4) [[C2GEN]])
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 noundef 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)* noundef 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]]
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 noundef 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)* noundef 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]]
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
139 __kernel void test##AS() { \
149 // COMMON-LABEL: @test__local
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)*))
154 // Test the address space of 'this' when invoking a method.
155 // COMMON: call spir_func noundef i32 @_ZNU3AS41C3getEv(%class.C addrspace(4)* {{[^,]*}} addrspacecast (%class.C addrspace(3)* @_ZZ11test__localE1c to %class.C addrspace(4)*))
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)* noundef 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)
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]])
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 noundef align 4 dereferenceable(4) %class.C addrspace(4)* @_ZNU3AS41CaSERU3AS4KS_(%class.C addrspace(4)* {{[^,]*}} [[C2GEN]], %class.C addrspace(4)* noundef 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]]
177 // COMMON-LABEL: @test__private
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]])
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 noundef i32 @_ZNU3AS41C3getEv(%class.C addrspace(4)* {{[^,]*}} [[CGEN]])
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)* noundef 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]]
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]])
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 noundef align 4 dereferenceable(4) %class.C addrspace(4)* @_ZNU3AS41CaSERU3AS4KS_(%class.C addrspace(4)* {{[^,]*}} [[C2GEN]], %class.C addrspace(4)* noundef 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]]
207 // Test that calling a const method from a non-const method does not crash Clang.
208 class ConstAndNonConstMethod {
210 void DoConst() const {