| 1 | // RUN: %clang_cc1 %s -triple spir-unknown-unknown -cl-std=c++ -emit-llvm -pedantic -verify -O0 -o - -DDECL | FileCheck %s --check-prefixes="COMMON,EXPL" |
| 2 | // RUN: %clang_cc1 %s -triple spir-unknown-unknown -cl-std=c++ -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=c++ -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 | |
| 14 | class C { |
| 15 | public: |
| 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) |
| 35 | C::C() { v = 2; }; |
| 36 | |
| 37 | C::C(C &&c) { v = c.v; } |
| 38 | |
| 39 | C::C(const C &c) { v = c.v; } |
| 40 | |
| 41 | C &C::operator=(const C &c) { |
| 42 | v = c.v; |
| 43 | return *this; |
| 44 | } |
| 45 | |
| 46 | C &C::operator=(C &&c) & { |
| 47 | v = c.v; |
| 48 | return *this; |
| 49 | } |
| 50 | #endif |
| 51 | |
| 52 | int C::outside() { |
| 53 | return v; |
| 54 | } |
| 55 | |
| 56 | extern C&& foo(); |
| 57 | |
| 58 | __global C c; |
| 59 | |
| 60 | __kernel void test__global() { |
| 61 | int i = c.get(); |
| 62 | int i2 = c.outside(); |
| 63 | C c1(c); |
| 64 | C c2; |
| 65 | c2 = c1; |
| 66 | C c3 = c1 + c2; |
| 67 | C c4(foo()); |
| 68 | C c5 = foo(); |
| 69 | } |
| 70 | |
| 71 | // Test that the address space is __generic for all members |
| 72 | // EXPL: @_ZNU3AS41CC2Ev(%class.C addrspace(4)* %this) |
| 73 | // EXPL: @_ZNU3AS41CC1Ev(%class.C addrspace(4)* %this) |
| 74 | // EXPL: @_ZNU3AS41CC2EOU3AS4S_(%class.C addrspace(4)* %this |
| 75 | // EXPL: @_ZNU3AS41CC1EOU3AS4S_(%class.C addrspace(4)* %this |
| 76 | // EXPL: @_ZNU3AS41CC2ERU3AS4KS_(%class.C addrspace(4)* %this |
| 77 | // EXPL: @_ZNU3AS41CC1ERU3AS4KS_(%class.C addrspace(4)* %this |
| 78 | // EXPL: @_ZNU3AS41CaSERU3AS4KS_(%class.C addrspace(4)* %this |
| 79 | // EXPL: @_ZNU3AS4R1CaSEOU3AS4S_(%class.C addrspace(4)* %this |
| 80 | // COMMON: @_ZNU3AS41C7outsideEv(%class.C addrspace(4)* %this) |
| 81 | |
| 82 | // EXPL-LABEL: @__cxx_global_var_init() |
| 83 | // EXPL: call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* addrspacecast (%class.C addrspace(1)* @c to %class.C addrspace(4)*)) |
| 84 | |
| 85 | // COMMON-LABEL: @_Z12test__globalv() |
| 86 | |
| 87 | // Test the address space of 'this' when invoking a method. |
| 88 | // COMMON: %call = call i32 @_ZNU3AS41C3getEv(%class.C addrspace(4)* addrspacecast (%class.C addrspace(1)* @c to %class.C addrspace(4)*)) |
| 89 | |
| 90 | // Test the address space of 'this' when invoking a method that is declared in the file contex. |
| 91 | // COMMON: %call1 = call i32 @_ZNU3AS41C7outsideEv(%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 copy-constructor. |
| 94 | // COMMON: [[C1GEN:%[0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)* |
| 95 | // IMPL: [[C1VOID:%[0-9]+]] = bitcast %class.C* %c1 to i8* |
| 96 | // 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)*) |
| 97 | // EXPL: call void @_ZNU3AS41CC1ERU3AS4KS_(%class.C addrspace(4)* [[C1GEN]], %class.C addrspace(4)* dereferenceable(4) addrspacecast (%class.C addrspace(1)* @c to %class.C addrspace(4)*)) |
| 98 | |
| 99 | // Test the address space of 'this' when invoking a constructor. |
| 100 | // EXPL: [[C2GEN:%[0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)* |
| 101 | // EXPL: call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* [[C2GEN]]) |
| 102 | |
| 103 | // Test the address space of 'this' when invoking assignment operator. |
| 104 | // COMMON: [[C1GEN:%[0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)* |
| 105 | // COMMON: [[C2GEN:%[0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)* |
| 106 | // EXPL: call dereferenceable(4) %class.C addrspace(4)* @_ZNU3AS41CaSERU3AS4KS_(%class.C addrspace(4)* [[C2GEN]], %class.C addrspace(4)* dereferenceable(4) [[C1GEN]]) |
| 107 | // IMPL: [[C2GENVOID:%[0-9]+]] = bitcast %class.C addrspace(4)* [[C2GEN]] to i8 addrspace(4)* |
| 108 | // IMPL: [[C1GENVOID:%[0-9]+]] = bitcast %class.C addrspace(4)* [[C1GEN]] to i8 addrspace(4)* |
| 109 | // IMPL: call void @llvm.memcpy.p4i8.p4i8.i32(i8 addrspace(4)* {{.*}}[[C2GENVOID]], i8 addrspace(4)* {{.*}}[[C1GENVOID]] |
| 110 | |
| 111 | // Test the address space of 'this' when invoking the operator+ |
| 112 | // COMMON: [[C3GEN:%[0-9]+]] = addrspacecast %class.C* %c3 to %class.C addrspace(4)* |
| 113 | // COMMON: [[C1GEN:%[0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)* |
| 114 | // COMMON: [[C2GEN:%[0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)* |
| 115 | // COMMON: call void @_ZNU3AS41CplERU3AS4KS_(%class.C* sret %ref.tmp, %class.C addrspace(4)* [[C1GEN]], %class.C addrspace(4)* dereferenceable(4) [[C2GEN]]) |
| 116 | // COMMON: [[REFGEN:%[0-9]+]] = addrspacecast %class.C* %ref.tmp to %class.C addrspace(4)* |
| 117 | // EXPL: call void @_ZNU3AS41CC1EOU3AS4S_(%class.C addrspace(4)* [[C3GEN]], %class.C addrspace(4)* dereferenceable(4) [[REFGEN]]) |
| 118 | // IMPL: [[C3VOID:%[0-9]+]] = bitcast %class.C* %c3 to i8* |
| 119 | // IMPL: [[REFGENVOID:%[0-9]+]] = bitcast %class.C addrspace(4)* [[REFGEN]] to i8 addrspace(4)* |
| 120 | // IMPL: call void @llvm.memcpy.p0i8.p4i8.i32(i8* {{.*}}[[C3VOID]], i8 addrspace(4)* {{.*}}[[REFGENVOID]] |
| 121 | |
| 122 | // Test the address space of 'this' when invoking the move constructor |
| 123 | // COMMON: [[C4GEN:%[0-9]+]] = addrspacecast %class.C* %c4 to %class.C addrspace(4)* |
| 124 | // COMMON: [[CALL:%call[0-9]+]] = call spir_func dereferenceable(4) %class.C addrspace(4)* @_Z3foov() |
| 125 | // EXPL: call void @_ZNU3AS41CC1EOU3AS4S_(%class.C addrspace(4)* [[C4GEN]], %class.C addrspace(4)* dereferenceable(4) [[CALL]]) |
| 126 | // IMPL: [[C4VOID:%[0-9]+]] = bitcast %class.C* %c4 to i8* |
| 127 | // IMPL: [[CALLVOID:%[0-9]+]] = bitcast %class.C addrspace(4)* [[CALL]] to i8 addrspace(4)* |
| 128 | // IMPL: call void @llvm.memcpy.p0i8.p4i8.i32(i8* {{.*}}[[C4VOID]], i8 addrspace(4)* {{.*}}[[CALLVOID]] |
| 129 | |
| 130 | // Test the address space of 'this' when invoking the move assignment |
| 131 | // COMMON: [[C5GEN:%[0-9]+]] = addrspacecast %class.C* %c5 to %class.C addrspace(4)* |
| 132 | // COMMON: [[CALL:%call[0-9]+]] = call spir_func dereferenceable(4) %class.C addrspace(4)* @_Z3foov() |
| 133 | // EXPL: call void @_ZNU3AS41CC1EOU3AS4S_(%class.C addrspace(4)* [[C5GEN:%[0-9]+]], %class.C addrspace(4)* dereferenceable(4) %call4) |
| 134 | // IMPL: [[C5VOID:%[0-9]+]] = bitcast %class.C* %c5 to i8* |
| 135 | // IMPL: [[CALLVOID:%[0-9]+]] = bitcast %class.C addrspace(4)* [[CALL]] to i8 addrspace(4)* |
| 136 | // IMPL: call void @llvm.memcpy.p0i8.p4i8.i32(i8* {{.*}}[[C5VOID]], i8 addrspace(4)* {{.*}}[[CALLVOID]] |
| 137 | |
| 138 | // Tests address space of inline members |
| 139 | //COMMON: @_ZNU3AS41C3getEv(%class.C addrspace(4)* %this) |
| 140 | //COMMON: @_ZNU3AS41CplERU3AS4KS_(%class.C* noalias sret %agg.result, %class.C addrspace(4)* %this |
| 141 | #define TEST(AS) \ |
| 142 | __kernel void test##AS() { \ |
| 143 | AS C c; \ |
| 144 | int i = c.get(); \ |
| 145 | C c1(c); \ |
| 146 | C c2; \ |
| 147 | c2 = c1; \ |
| 148 | } |
| 149 | |
| 150 | TEST(__local) |
| 151 | |
| 152 | // COMMON-LABEL: _Z11test__localv |
| 153 | // EXPL: @__cxa_guard_acquire |
| 154 | |
| 155 | // Test the address space of 'this' when invoking a constructor for an object in non-default address space |
| 156 | // EXPL: call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* addrspacecast (%class.C addrspace(3)* @_ZZ11test__localvE1c to %class.C addrspace(4)*)) |
| 157 | |
| 158 | // Test the address space of 'this' when invoking a method. |
| 159 | // COMMON: %call = call i32 @_ZNU3AS41C3getEv(%class.C addrspace(4)* addrspacecast (%class.C addrspace(3)* @_ZZ11test__localvE1c to %class.C addrspace(4)*)) |
| 160 | |
| 161 | |
| 162 | // Test the address space of 'this' when invoking copy-constructor. |
| 163 | // COMMON: [[C1GEN:%[0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)* |
| 164 | // EXPL: call void @_ZNU3AS41CC1ERU3AS4KS_(%class.C addrspace(4)* [[C1GEN]], %class.C addrspace(4)* dereferenceable(4) addrspacecast (%class.C addrspace(3)* @_ZZ11test__localvE1c to %class.C addrspace(4)*)) |
| 165 | // IMPL: [[C1VOID:%[0-9]+]] = bitcast %class.C* %c1 to i8* |
| 166 | // IMPL: call void @llvm.memcpy.p0i8.p4i8.i32(i8* {{.*}}[[C1VOID]], i8 addrspace(4)* {{.*}}addrspacecast (i8 addrspace(3)* bitcast (%class.C addrspace(3)* @_ZZ11test__localvE1c to i8 addrspace(3)*) to i8 addrspace(4)*), i32 4, i1 false) |
| 167 | |
| 168 | // Test the address space of 'this' when invoking a constructor. |
| 169 | // EXPL: [[C2GEN:%[0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)* |
| 170 | // EXPL: call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* [[C2GEN]]) |
| 171 | |
| 172 | // Test the address space of 'this' when invoking assignment operator. |
| 173 | // COMMON: [[C1GEN:%[0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)* |
| 174 | // COMMON: [[C2GEN:%[0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)* |
| 175 | // EXPL: call dereferenceable(4) %class.C addrspace(4)* @_ZNU3AS41CaSERU3AS4KS_(%class.C addrspace(4)* [[C2GEN]], %class.C addrspace(4)* dereferenceable(4) [[C1GEN]]) |
| 176 | // IMPL: [[C2GENVOID:%[0-9]+]] = bitcast %class.C addrspace(4)* [[C2GEN]] to i8 addrspace(4)* |
| 177 | // IMPL: [[C1GENVOID:%[0-9]+]] = bitcast %class.C addrspace(4)* [[C1GEN]] to i8 addrspace(4)* |
| 178 | // IMPL: call void @llvm.memcpy.p4i8.p4i8.i32(i8 addrspace(4)* {{.*}}[[C2GENVOID]], i8 addrspace(4)* {{.*}}[[C1GENVOID]] |
| 179 | |
| 180 | TEST(__private) |
| 181 | |
| 182 | // CHECK-LABEL: @_Z13test__privatev |
| 183 | |
| 184 | // Test the address space of 'this' when invoking a constructor for an object in non-default address space |
| 185 | // EXPL: [[CGEN:%[0-9]+]] = addrspacecast %class.C* %c to %class.C addrspace(4)* |
| 186 | // EXPL: call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* [[CGEN]]) |
| 187 | |
| 188 | // Test the address space of 'this' when invoking a method. |
| 189 | // COMMON: [[CGEN:%[0-9]+]] = addrspacecast %class.C* %c to %class.C addrspace(4)* |
| 190 | // COMMON: %call = call i32 @_ZNU3AS41C3getEv(%class.C addrspace(4)* [[CGEN]]) |
| 191 | |
| 192 | // Test the address space of 'this' when invoking a copy-constructor. |
| 193 | // COMMON: [[C1GEN:%[0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)* |
| 194 | // COMMON: [[CGEN:%[0-9]+]] = addrspacecast %class.C* %c to %class.C addrspace(4)* |
| 195 | // EXPL: call void @_ZNU3AS41CC1ERU3AS4KS_(%class.C addrspace(4)* [[C1GEN]], %class.C addrspace(4)* dereferenceable(4) [[CGEN]]) |
| 196 | // IMPL: [[C1VOID:%[0-9]+]] = bitcast %class.C* %c1 to i8* |
| 197 | // IMPL: [[CGENVOID:%[0-9]+]] = bitcast %class.C addrspace(4)* [[CGEN]] to i8 addrspace(4)* |
| 198 | // IMPL: call void @llvm.memcpy.p0i8.p4i8.i32(i8* {{.*}}[[C1VOID]], i8 addrspace(4)* {{.*}}[[CGENVOID]] |
| 199 | |
| 200 | // Test the address space of 'this' when invoking a constructor. |
| 201 | // EXPL: [[C2GEN:%[0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)* |
| 202 | // EXPL: call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* [[C2GEN]]) |
| 203 | |
| 204 | // Test the address space of 'this' when invoking a copy-assignment. |
| 205 | // COMMON: [[C1GEN:%[0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)* |
| 206 | // COMMON: [[C2GEN:%[0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)* |
| 207 | // EXPL: %call1 = call dereferenceable(4) %class.C addrspace(4)* @_ZNU3AS41CaSERU3AS4KS_(%class.C addrspace(4)* [[C2GEN]], %class.C addrspace(4)* dereferenceable(4) [[C1GEN]]) |
| 208 | // IMPL: [[C2GENVOID:%[0-9]+]] = bitcast %class.C addrspace(4)* [[C2GEN]] to i8 addrspace(4)* |
| 209 | // IMPL: [[C1GENVOID:%[0-9]+]] = bitcast %class.C addrspace(4)* [[C1GEN]] to i8 addrspace(4)* |
| 210 | // IMPL: call void @llvm.memcpy.p4i8.p4i8.i32(i8 addrspace(4)* {{.*}}[[C2GENVOID]], i8 addrspace(4)* {{.*}}[[C1GENVOID]] |
| 211 | |