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