204 lines · plain
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-diagnostics5 6// Test that the 'this' pointer is in the __generic address space.7 8#ifdef USE_DEFLT9#define DEFAULT =default10#else11#define DEFAULT12#endif13 14class C {15public:16 int v;17#ifdef DECL18 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#endif24 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#endif51 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 members73// EXPL: @_ZNU3AS41CC2Ev(ptr addrspace(4) {{[^,]*}} %this)74// EXPL: @_ZNU3AS41CC1Ev(ptr addrspace(4) {{[^,]*}} %this)75// EXPL: @_ZNU3AS41CC2EOU3AS4S_(ptr addrspace(4) {{[^,]*}} %this76// EXPL: @_ZNU3AS41CC1EOU3AS4S_(ptr addrspace(4) {{[^,]*}} %this77// EXPL: @_ZNU3AS41CC2ERU3AS4KS_(ptr addrspace(4) {{[^,]*}} %this78// EXPL: @_ZNU3AS41CC1ERU3AS4KS_(ptr addrspace(4) {{[^,]*}} %this79// EXPL: @_ZNU3AS41CaSERU3AS4KS_(ptr addrspace(4) {{[^,]*}} %this80// EXPL: @_ZNU3AS4R1CaSEOU3AS4S_(ptr addrspace(4) {{[^,]*}} %this81// 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 constructor117// 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 assignment123// 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#define TEST(AS) \129 __kernel void test##AS() { \130 AS C c; \131 int i = c.get(); \132 C c1(c); \133 C c2; \134 c2 = c1; \135 }136 137TEST(__local)138 139// COMMON-LABEL: @test__local140 141// Test that we don't initialize an object in local address space.142// EXPL-NOT: call spir_func void @_ZNU3AS41CC1Ev(ptr addrspace(4) addrspacecast (ptr addrspace(3) @_ZZ11test__localE1c to ptr addrspace(4)))143 144// Test the address space of 'this' when invoking a method.145// COMMON: call spir_func noundef i32 @_ZNU3AS41C3getEv(ptr addrspace(4) {{[^,]*}} addrspacecast (ptr addrspace(3) @_ZZ11test__localE1c to ptr addrspace(4)))146 147// Test the address space of 'this' when invoking copy-constructor.148// COMMON: [[C1GEN:%[.a-z0-9]+]] = addrspacecast ptr %c1 to ptr addrspace(4)149// 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)))150// 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)151 152// Test the address space of 'this' when invoking a constructor.153// EXPL: [[C2GEN:%[.a-z0-9]+]] = addrspacecast ptr %c2 to ptr addrspace(4)154// EXPL: call spir_func void @_ZNU3AS41CC1Ev(ptr addrspace(4) {{[^,]*}} [[C2GEN]])155 156// Test the address space of 'this' when invoking assignment operator.157// COMMON: [[C1GEN:%[.a-z0-9]+]] = addrspacecast ptr %c1 to ptr addrspace(4)158// COMMON: [[C2GEN:%[.a-z0-9]+]] = addrspacecast ptr %c2 to ptr addrspace(4)159// 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]])160// IMPL: call void @llvm.memcpy.p4.p4.i32(ptr addrspace(4) {{.*}}[[C2GEN]], ptr addrspace(4) {{.*}}[[C1GEN]]161 162TEST(__private)163 164// COMMON-LABEL: @test__private165 166// Test the address space of 'this' when invoking a constructor for an object in non-default address space167// EXPL: [[CGEN:%[.a-z0-9]+]] = addrspacecast ptr %c to ptr addrspace(4)168// EXPL: call spir_func void @_ZNU3AS41CC1Ev(ptr addrspace(4) {{[^,]*}} [[CGEN]])169 170// Test the address space of 'this' when invoking a method.171// COMMON: [[CGEN:%[.a-z0-9]+]] = addrspacecast ptr %c to ptr addrspace(4)172// COMMON: call spir_func noundef i32 @_ZNU3AS41C3getEv(ptr addrspace(4) {{[^,]*}} [[CGEN]])173 174// Test the address space of 'this' when invoking a copy-constructor.175// COMMON: [[C1GEN:%[.a-z0-9]+]] = addrspacecast ptr %c1 to ptr addrspace(4)176// COMMON: [[CGEN:%[.a-z0-9]+]] = addrspacecast ptr %c to ptr addrspace(4)177// EXPL: call spir_func void @_ZNU3AS41CC1ERU3AS4KS_(ptr addrspace(4) {{[^,]*}} [[C1GEN]], ptr addrspace(4) noundef align 4 dereferenceable(4) [[CGEN]])178// IMPL: call void @llvm.memcpy.p0.p4.i32(ptr {{.*}}%c1, ptr addrspace(4) {{.*}}[[CGEN]]179 180// Test the address space of 'this' when invoking a constructor.181// EXPL: [[C2GEN:%[.a-z0-9]+]] = addrspacecast ptr %c2 to ptr addrspace(4)182// EXPL: call spir_func void @_ZNU3AS41CC1Ev(ptr addrspace(4) {{[^,]*}} [[C2GEN]])183 184// Test the address space of 'this' when invoking a copy-assignment.185// COMMON: [[C1GEN:%[.a-z0-9]+]] = addrspacecast ptr %c1 to ptr addrspace(4)186// COMMON: [[C2GEN:%[.a-z0-9]+]] = addrspacecast ptr %c2 to ptr addrspace(4)187// 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]])188// IMPL: call void @llvm.memcpy.p4.p4.i32(ptr addrspace(4) {{.*}}[[C2GEN]], ptr addrspace(4) {{.*}}[[C1GEN]]189 190// Tests address space of inline members191//COMMON: @_ZNU3AS41C3getEv(ptr addrspace(4) {{[^,]*}} %this)192//COMMON: @_ZNU3AS41CplERU3AS4KS_(ptr dead_on_unwind noalias writable sret(%class.C) align 4 %agg.result, ptr addrspace(4) {{[^,]*}} %this193 194// Test that calling a const method from a non-const method does not crash Clang.195class ConstAndNonConstMethod {196public:197 void DoConst() const {198 }199 200 void DoNonConst() {201 DoConst();202 }203};204