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