diff options
Diffstat (limited to 'test/CodeGenOpenCLCXX')
| -rw-r--r-- | test/CodeGenOpenCLCXX/address-space-castoperators.cpp | 14 | ||||
| -rw-r--r-- | test/CodeGenOpenCLCXX/addrspace-derived-base.cl | 22 | ||||
| -rw-r--r-- | test/CodeGenOpenCLCXX/addrspace-of-this.cl | 223 | ||||
| -rw-r--r-- | test/CodeGenOpenCLCXX/addrspace-operators.cl | 53 | ||||
| -rw-r--r-- | test/CodeGenOpenCLCXX/addrspace-references.cl | 14 | ||||
| -rw-r--r-- | test/CodeGenOpenCLCXX/local_addrspace_init.cl | 20 | ||||
| -rw-r--r-- | test/CodeGenOpenCLCXX/method-overload-address-space.cl | 35 |
7 files changed, 281 insertions, 100 deletions
diff --git a/test/CodeGenOpenCLCXX/address-space-castoperators.cpp b/test/CodeGenOpenCLCXX/address-space-castoperators.cpp new file mode 100644 index 0000000000..e5c3ee880c --- /dev/null +++ b/test/CodeGenOpenCLCXX/address-space-castoperators.cpp @@ -0,0 +1,14 @@ +// RUN: %clang_cc1 %s -triple spir-unknown-unknown -cl-std=c++ -emit-llvm -O0 -o - | FileCheck %s + +void test_reinterpret_cast(){ +__private float x; +__private float& y = x; +// We don't need bitcast to cast pointer type and +// address space at the same time. +//CHECK: addrspacecast float* %x to i32 addrspace(4)* +//CHECK: [[REG:%[0-9]+]] = load float*, float** %y +//CHECK: addrspacecast float* [[REG]] to i32 addrspace(4)* +//CHECK-NOT: bitcast +__generic int& rc1 = reinterpret_cast<__generic int&>(x); +__generic int& rc2 = reinterpret_cast<__generic int&>(y); +} diff --git a/test/CodeGenOpenCLCXX/addrspace-derived-base.cl b/test/CodeGenOpenCLCXX/addrspace-derived-base.cl new file mode 100644 index 0000000000..59e29ee417 --- /dev/null +++ b/test/CodeGenOpenCLCXX/addrspace-derived-base.cl @@ -0,0 +1,22 @@ +// RUN: %clang_cc1 %s -triple spir -cl-std=c++ -emit-llvm -O0 -o - | FileCheck %s + +struct B { + int mb; +}; + +class D : public B { +public: + int getmb() { return mb; } +}; + +void foo() { + D d; + //CHECK: addrspacecast %class.D* %d to %class.D addrspace(4)* + //CHECK: call i32 @_ZNU3AS41D5getmbEv(%class.D addrspace(4)* + d.getmb(); +} + +//Derived and Base are in the same address space. + +//CHECK: define linkonce_odr i32 @_ZNU3AS41D5getmbEv(%class.D addrspace(4)* %this) +//CHECK: bitcast %class.D addrspace(4)* %this1 to %struct.B addrspace(4)* diff --git a/test/CodeGenOpenCLCXX/addrspace-of-this.cl b/test/CodeGenOpenCLCXX/addrspace-of-this.cl index 83af3fbcb3..fee804e161 100644 --- a/test/CodeGenOpenCLCXX/addrspace-of-this.cl +++ b/test/CodeGenOpenCLCXX/addrspace-of-this.cl @@ -1,25 +1,26 @@ -// RUN: %clang_cc1 %s -triple spir-unknown-unknown -cl-std=c++ -emit-llvm -pedantic -verify -O0 -o - | FileCheck %s +// RUN: %clang_cc1 %s -triple spir-unknown-unknown -cl-std=c++ -emit-llvm -pedantic -verify -O0 -o - -DDECL | FileCheck %s --check-prefixes="COMMON,EXPL" +// 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" +// RUN: %clang_cc1 %s -triple spir-unknown-unknown -cl-std=c++ -emit-llvm -pedantic -verify -O0 -o - | FileCheck %s --check-prefixes="COMMON,IMPL" // expected-no-diagnostics // Test that the 'this' pointer is in the __generic address space. -// FIXME: Add support for __constant address space. +#ifdef USE_DEFLT +#define DEFAULT =default +#else +#define DEFAULT +#endif class C { public: int v; - C() { v = 2; } - C(C &&c) { v = c.v; } - C(const C &c) { v = c.v; } - C &operator=(const C &c) { - v = c.v; - return *this; - } - C &operator=(C &&c) & { - v = c.v; - return *this; - } - +#ifdef DECL + C() DEFAULT; + C(C &&c) DEFAULT; + C(const C &c) DEFAULT; + C &operator=(const C &c) DEFAULT; + C &operator=(C &&c) & DEFAULT; +#endif C operator+(const C& c) { v += c.v; return *this; @@ -30,6 +31,24 @@ public: int outside(); }; +#if defined(DECL) && !defined(USE_DEFLT) +C::C() { v = 2; }; + +C::C(C &&c) { v = c.v; } + +C::C(const C &c) { v = c.v; } + +C &C::operator=(const C &c) { + v = c.v; + return *this; +} + +C &C::operator=(C &&c) & { + v = c.v; + return *this; +} +#endif + int C::outside() { return v; } @@ -40,7 +59,8 @@ __global C c; __kernel void test__global() { int i = c.get(); - int i2 = c.outside(); + int i2 = (&c)->get(); + int i3 = c.outside(); C c1(c); C c2; c2 = c1; @@ -49,58 +69,78 @@ __kernel void test__global() { C c5 = foo(); } -// CHECK-LABEL: @__cxx_global_var_init() -// CHECK: call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* addrspacecast (%class.C addrspace(1)* @c to %class.C addrspace(4)*)) +// Test that the address space is __generic for all members +// EXPL: @_ZNU3AS41CC2Ev(%class.C addrspace(4)* %this) +// EXPL: @_ZNU3AS41CC1Ev(%class.C addrspace(4)* %this) +// EXPL: @_ZNU3AS41CC2EOU3AS4S_(%class.C addrspace(4)* %this +// EXPL: @_ZNU3AS41CC1EOU3AS4S_(%class.C addrspace(4)* %this +// EXPL: @_ZNU3AS41CC2ERU3AS4KS_(%class.C addrspace(4)* %this +// EXPL: @_ZNU3AS41CC1ERU3AS4KS_(%class.C addrspace(4)* %this +// EXPL: @_ZNU3AS41CaSERU3AS4KS_(%class.C addrspace(4)* %this +// EXPL: @_ZNU3AS4R1CaSEOU3AS4S_(%class.C addrspace(4)* %this +// COMMON: @_ZNU3AS41C7outsideEv(%class.C addrspace(4)* %this) -// Test that the address space is __generic for the constructor -// CHECK-LABEL: @_ZNU3AS41CC1Ev(%class.C addrspace(4)* %this) -// CHECK: entry: -// CHECK: %this.addr = alloca %class.C addrspace(4)*, align 4 -// CHECK: store %class.C addrspace(4)* %this, %class.C addrspace(4)** %this.addr, align 4 -// CHECK: %this1 = load %class.C addrspace(4)*, %class.C addrspace(4)** %this.addr, align 4 -// CHECK: call void @_ZNU3AS41CC2Ev(%class.C addrspace(4)* %this1) -// CHECK: ret void +// EXPL-LABEL: @__cxx_global_var_init() +// EXPL: call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* addrspacecast (%class.C addrspace(1)* @c to %class.C addrspace(4)*)) -// CHECK-LABEL: @_Z12test__globalv() +// COMMON-LABEL: @_Z12test__globalv() // Test the address space of 'this' when invoking a method. -// CHECK: %call = call i32 @_ZNU3AS41C3getEv(%class.C addrspace(4)* addrspacecast (%class.C addrspace(1)* @c to %class.C addrspace(4)*)) +// COMMON: call i32 @_ZNU3AS41C3getEv(%class.C addrspace(4)* addrspacecast (%class.C addrspace(1)* @c to %class.C addrspace(4)*)) +// Test the address space of 'this' when invoking a method using a pointer to the object. +// COMMON: call i32 @_ZNU3AS41C3getEv(%class.C addrspace(4)* addrspacecast (%class.C addrspace(1)* @c to %class.C addrspace(4)*)) // Test the address space of 'this' when invoking a method that is declared in the file contex. -// CHECK: %call1 = call i32 @_ZNU3AS41C7outsideEv(%class.C addrspace(4)* addrspacecast (%class.C addrspace(1)* @c to %class.C addrspace(4)*)) +// COMMON: call i32 @_ZNU3AS41C7outsideEv(%class.C addrspace(4)* addrspacecast (%class.C addrspace(1)* @c to %class.C addrspace(4)*)) // Test the address space of 'this' when invoking copy-constructor. -// CHECK: [[C1GEN:%[0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)* -// CHECK: 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)*)) +// COMMON: [[C1GEN:%[0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)* +// IMPL: [[C1VOID:%[0-9]+]] = bitcast %class.C* %c1 to i8* +// 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)*) +// 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)*)) // Test the address space of 'this' when invoking a constructor. -// CHECK: [[C2GEN:%[0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)* -// CHECK: call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* [[C2GEN]]) +// EXPL: [[C2GEN:%[0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)* +// EXPL: call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* [[C2GEN]]) // Test the address space of 'this' when invoking assignment operator. -// CHECK: [[C1GEN:%[0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)* -// CHECK: [[C2GEN:%[0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)* -// CHECK: %call2 = call dereferenceable(4) %class.C addrspace(4)* @_ZNU3AS41CaSERU3AS4KS_(%class.C addrspace(4)* [[C2GEN]], %class.C addrspace(4)* dereferenceable(4) [[C1GEN]]) +// COMMON: [[C1GEN:%[0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)* +// COMMON: [[C2GEN:%[0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)* +// EXPL: call dereferenceable(4) %class.C addrspace(4)* @_ZNU3AS41CaSERU3AS4KS_(%class.C addrspace(4)* [[C2GEN]], %class.C addrspace(4)* dereferenceable(4) [[C1GEN]]) +// IMPL: [[C2GENVOID:%[0-9]+]] = bitcast %class.C addrspace(4)* [[C2GEN]] to i8 addrspace(4)* +// IMPL: [[C1GENVOID:%[0-9]+]] = bitcast %class.C addrspace(4)* [[C1GEN]] to i8 addrspace(4)* +// IMPL: call void @llvm.memcpy.p4i8.p4i8.i32(i8 addrspace(4)* {{.*}}[[C2GENVOID]], i8 addrspace(4)* {{.*}}[[C1GENVOID]] // Test the address space of 'this' when invoking the operator+ -// CHECK: [[C3GEN:%[0-9]+]] = addrspacecast %class.C* %c3 to %class.C addrspace(4)* -// CHECK: [[C1GEN:%[0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)* -// CHECK: [[C2GEN:%[0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)* -// CHECK: call void @_ZNU3AS41CplERU3AS4KS_(%class.C* sret %ref.tmp, %class.C addrspace(4)* [[C1GEN]], %class.C addrspace(4)* dereferenceable(4) [[C2GEN]]) -// CHECK: [[REFGEN:%[0-9]+]] = addrspacecast %class.C* %ref.tmp to %class.C addrspace(4)* -// CHECK: call void @_ZNU3AS41CC1EOU3AS4S_(%class.C addrspace(4)* [[C3GEN]], %class.C addrspace(4)* dereferenceable(4) [[REFGEN]]) +// COMMON: [[C3GEN:%[0-9]+]] = addrspacecast %class.C* %c3 to %class.C addrspace(4)* +// COMMON: [[C1GEN:%[0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)* +// COMMON: [[C2GEN:%[0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)* +// COMMON: call void @_ZNU3AS41CplERU3AS4KS_(%class.C* sret %ref.tmp, %class.C addrspace(4)* [[C1GEN]], %class.C addrspace(4)* dereferenceable(4) [[C2GEN]]) +// COMMON: [[REFGEN:%[0-9]+]] = addrspacecast %class.C* %ref.tmp to %class.C addrspace(4)* +// EXPL: call void @_ZNU3AS41CC1EOU3AS4S_(%class.C addrspace(4)* [[C3GEN]], %class.C addrspace(4)* dereferenceable(4) [[REFGEN]]) +// IMPL: [[C3VOID:%[0-9]+]] = bitcast %class.C* %c3 to i8* +// IMPL: [[REFGENVOID:%[0-9]+]] = bitcast %class.C addrspace(4)* [[REFGEN]] to i8 addrspace(4)* +// IMPL: call void @llvm.memcpy.p0i8.p4i8.i32(i8* {{.*}}[[C3VOID]], i8 addrspace(4)* {{.*}}[[REFGENVOID]] // Test the address space of 'this' when invoking the move constructor -// CHECK: [[C4GEN:%[0-9]+]] = addrspacecast %class.C* %c4 to %class.C addrspace(4)* -// CHECK: %call3 = call spir_func dereferenceable(4) %class.C addrspace(4)* @_Z3foov() -// CHECK: call void @_ZNU3AS41CC1EOU3AS4S_(%class.C addrspace(4)* [[C4GEN]], %class.C addrspace(4)* dereferenceable(4) %call3) +// COMMON: [[C4GEN:%[0-9]+]] = addrspacecast %class.C* %c4 to %class.C addrspace(4)* +// COMMON: [[CALL:%call[0-9]+]] = call spir_func dereferenceable(4) %class.C addrspace(4)* @_Z3foov() +// EXPL: call void @_ZNU3AS41CC1EOU3AS4S_(%class.C addrspace(4)* [[C4GEN]], %class.C addrspace(4)* dereferenceable(4) [[CALL]]) +// IMPL: [[C4VOID:%[0-9]+]] = bitcast %class.C* %c4 to i8* +// IMPL: [[CALLVOID:%[0-9]+]] = bitcast %class.C addrspace(4)* [[CALL]] to i8 addrspace(4)* +// IMPL: call void @llvm.memcpy.p0i8.p4i8.i32(i8* {{.*}}[[C4VOID]], i8 addrspace(4)* {{.*}}[[CALLVOID]] // Test the address space of 'this' when invoking the move assignment -// CHECK: [[C5GEN:%[0-9]+]] = addrspacecast %class.C* %c5 to %class.C addrspace(4)* -// CHECK: %call4 = call spir_func dereferenceable(4) %class.C addrspace(4)* @_Z3foov() #5 -// CHECK: call void @_ZNU3AS41CC1EOU3AS4S_(%class.C addrspace(4)* [[C5GEN:%[0-9]+]], %class.C addrspace(4)* dereferenceable(4) %call4) - - +// COMMON: [[C5GEN:%[0-9]+]] = addrspacecast %class.C* %c5 to %class.C addrspace(4)* +// COMMON: [[CALL:%call[0-9]+]] = call spir_func dereferenceable(4) %class.C addrspace(4)* @_Z3foov() +// EXPL: call void @_ZNU3AS41CC1EOU3AS4S_(%class.C addrspace(4)* [[C5GEN:%[0-9]+]], %class.C addrspace(4)* dereferenceable(4) [[CALL]]) +// IMPL: [[C5VOID:%[0-9]+]] = bitcast %class.C* %c5 to i8* +// IMPL: [[CALLVOID:%[0-9]+]] = bitcast %class.C addrspace(4)* [[CALL]] to i8 addrspace(4)* +// IMPL: call void @llvm.memcpy.p0i8.p4i8.i32(i8* {{.*}}[[C5VOID]], i8 addrspace(4)* {{.*}}[[CALLVOID]] + +// Tests address space of inline members +//COMMON: @_ZNU3AS41C3getEv(%class.C addrspace(4)* %this) +//COMMON: @_ZNU3AS41CplERU3AS4KS_(%class.C* noalias sret %agg.result, %class.C addrspace(4)* %this #define TEST(AS) \ __kernel void test##AS() { \ AS C c; \ @@ -112,77 +152,60 @@ __kernel void test__global() { TEST(__local) -// CHECK-LABEL: _Z11test__localv -// CHECK: @__cxa_guard_acquire +// COMMON-LABEL: _Z11test__localv -// Test the address space of 'this' when invoking a constructor for an object in non-default address space -// CHECK: call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* addrspacecast (%class.C addrspace(3)* @_ZZ11test__localvE1c to %class.C addrspace(4)*)) +// Test that we don't initialize an object in local address space. +// EXPL-NOT: call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* addrspacecast (%class.C addrspace(3)* @_ZZ11test__localvE1c to %class.C addrspace(4)*)) // Test the address space of 'this' when invoking a method. -// CHECK: %call = call i32 @_ZNU3AS41C3getEv(%class.C addrspace(4)* addrspacecast (%class.C addrspace(3)* @_ZZ11test__localvE1c to %class.C addrspace(4)*)) - +// COMMON: call i32 @_ZNU3AS41C3getEv(%class.C addrspace(4)* addrspacecast (%class.C addrspace(3)* @_ZZ11test__localvE1c to %class.C addrspace(4)*)) // Test the address space of 'this' when invoking copy-constructor. -// CHECK: [[C1GEN:%[0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)* -// CHECK: 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)*)) +// COMMON: [[C1GEN:%[0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)* +// 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)*)) +// IMPL: [[C1VOID:%[0-9]+]] = bitcast %class.C* %c1 to i8* +// 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) // Test the address space of 'this' when invoking a constructor. -// CHECK: [[C2GEN:%[0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)* -// CHECK: call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* [[C2GEN]]) +// EXPL: [[C2GEN:%[0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)* +// EXPL: call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* [[C2GEN]]) // Test the address space of 'this' when invoking assignment operator. -// CHECK: [[C1GEN:%[0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)* -// CHECK: [[C2GEN:%[0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)* -// CHECK: %call1 = call dereferenceable(4) %class.C addrspace(4)* @_ZNU3AS41CaSERU3AS4KS_(%class.C addrspace(4)* [[C2GEN]], %class.C addrspace(4)* dereferenceable(4) [[C1GEN]]) +// COMMON: [[C1GEN:%[0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)* +// COMMON: [[C2GEN:%[0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)* +// EXPL: call dereferenceable(4) %class.C addrspace(4)* @_ZNU3AS41CaSERU3AS4KS_(%class.C addrspace(4)* [[C2GEN]], %class.C addrspace(4)* dereferenceable(4) [[C1GEN]]) +// IMPL: [[C2GENVOID:%[0-9]+]] = bitcast %class.C addrspace(4)* [[C2GEN]] to i8 addrspace(4)* +// IMPL: [[C1GENVOID:%[0-9]+]] = bitcast %class.C addrspace(4)* [[C1GEN]] to i8 addrspace(4)* +// IMPL: call void @llvm.memcpy.p4i8.p4i8.i32(i8 addrspace(4)* {{.*}}[[C2GENVOID]], i8 addrspace(4)* {{.*}}[[C1GENVOID]] TEST(__private) // CHECK-LABEL: @_Z13test__privatev // Test the address space of 'this' when invoking a constructor for an object in non-default address space -// CHECK: [[CGEN:%[0-9]+]] = addrspacecast %class.C* %c to %class.C addrspace(4)* -// CHECK: call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* [[CGEN]]) - -// Test the address space of 'this' when invoking a method. -// CHECK: [[CGEN:%[0-9]+]] = addrspacecast %class.C* %c to %class.C addrspace(4)* -// CHECK: %call = call i32 @_ZNU3AS41C3getEv(%class.C addrspace(4)* [[CGEN]]) - -// Test the address space of 'this' when invoking a copy-constructor. -// CHECK: [[C1GEN:%[0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)* -// CHECK: [[CGEN:%[0-9]+]] = addrspacecast %class.C* %c to %class.C addrspace(4)* -// CHECK: call void @_ZNU3AS41CC1ERU3AS4KS_(%class.C addrspace(4)* [[C1GEN]], %class.C addrspace(4)* dereferenceable(4) [[CGEN]]) - -// Test the address space of 'this' when invoking a constructor. -// CHECK: [[C2GEN:%[0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)* -// CHECK: call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* [[C2GEN]]) - -// Test the address space of 'this' when invoking a copy-assignment. -// CHECK: [[C1GEN:%[0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)* -// CHECK: [[C2GEN:%[0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)* -// CHECK: %call1 = call dereferenceable(4) %class.C addrspace(4)* @_ZNU3AS41CaSERU3AS4KS_(%class.C addrspace(4)* [[C2GEN]], %class.C addrspace(4)* dereferenceable(4) [[C1GEN]]) - -TEST() - -// CHECK-LABEL: @_Z4testv() - -// Test the address space of 'this' when invoking a constructor for an object in non-default address space -// CHECK: [[CGEN:%[0-9]+]] = addrspacecast %class.C* %c to %class.C addrspace(4)* -// CHECK: call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* [[CGEN]]) +// EXPL: [[CGEN:%[0-9]+]] = addrspacecast %class.C* %c to %class.C addrspace(4)* +// EXPL: call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* [[CGEN]]) // Test the address space of 'this' when invoking a method. -// CHECK: [[CGEN:%[0-9]+]] = addrspacecast %class.C* %c to %class.C addrspace(4)* -// CHECK: %call = call i32 @_ZNU3AS41C3getEv(%class.C addrspace(4)* [[CGEN]]) +// COMMON: [[CGEN:%[0-9]+]] = addrspacecast %class.C* %c to %class.C addrspace(4)* +// COMMON: call i32 @_ZNU3AS41C3getEv(%class.C addrspace(4)* [[CGEN]]) // Test the address space of 'this' when invoking a copy-constructor. -// CHECK: [[C1GEN:%[0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)* -// CHECK: [[CGEN:%[0-9]+]] = addrspacecast %class.C* %c to %class.C addrspace(4)* -// CHECK: call void @_ZNU3AS41CC1ERU3AS4KS_(%class.C addrspace(4)* [[C1GEN]], %class.C addrspace(4)* dereferenceable(4) [[CGEN]]) +// COMMON: [[C1GEN:%[0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)* +// COMMON: [[CGEN:%[0-9]+]] = addrspacecast %class.C* %c to %class.C addrspace(4)* +// EXPL: call void @_ZNU3AS41CC1ERU3AS4KS_(%class.C addrspace(4)* [[C1GEN]], %class.C addrspace(4)* dereferenceable(4) [[CGEN]]) +// IMPL: [[C1VOID:%[0-9]+]] = bitcast %class.C* %c1 to i8* +// IMPL: [[CGENVOID:%[0-9]+]] = bitcast %class.C addrspace(4)* [[CGEN]] to i8 addrspace(4)* +// IMPL: call void @llvm.memcpy.p0i8.p4i8.i32(i8* {{.*}}[[C1VOID]], i8 addrspace(4)* {{.*}}[[CGENVOID]] // Test the address space of 'this' when invoking a constructor. -// CHECK: [[C2GEN:%[0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)* -// CHECK: call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* [[C2GEN]]) +// EXPL: [[C2GEN:%[0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)* +// EXPL: call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* [[C2GEN]]) // Test the address space of 'this' when invoking a copy-assignment. -// CHECK: [[C1GEN:%[0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)* -// CHECK: [[C2GEN:%[0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)* -// CHECK: %call1 = call dereferenceable(4) %class.C addrspace(4)* @_ZNU3AS41CaSERU3AS4KS_(%class.C addrspace(4)* [[C2GEN]], %class.C addrspace(4)* dereferenceable(4) [[C1GEN]]) +// COMMON: [[C1GEN:%[0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)* +// COMMON: [[C2GEN:%[0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)* +// EXPL: call dereferenceable(4) %class.C addrspace(4)* @_ZNU3AS41CaSERU3AS4KS_(%class.C addrspace(4)* [[C2GEN]], %class.C addrspace(4)* dereferenceable(4) [[C1GEN]]) +// IMPL: [[C2GENVOID:%[0-9]+]] = bitcast %class.C addrspace(4)* [[C2GEN]] to i8 addrspace(4)* +// IMPL: [[C1GENVOID:%[0-9]+]] = bitcast %class.C addrspace(4)* [[C1GEN]] to i8 addrspace(4)* +// IMPL: call void @llvm.memcpy.p4i8.p4i8.i32(i8 addrspace(4)* {{.*}}[[C2GENVOID]], i8 addrspace(4)* {{.*}}[[C1GENVOID]] diff --git a/test/CodeGenOpenCLCXX/addrspace-operators.cl b/test/CodeGenOpenCLCXX/addrspace-operators.cl new file mode 100644 index 0000000000..1f3389defd --- /dev/null +++ b/test/CodeGenOpenCLCXX/addrspace-operators.cl @@ -0,0 +1,53 @@ +//RUN: %clang_cc1 %s -triple spir -cl-std=c++ -emit-llvm -O0 -o - | FileCheck %s + +enum E { + a, + b, +}; + +class C { +public: + void Assign(E e) { me = e; } + void OrAssign(E e) { mi |= e; } + E me; + int mi; +}; + +__global E globE; +volatile __global int globVI; +__global int globI; +//CHECK-LABEL: define spir_func void @_Z3barv() +void bar() { + C c; + //CHECK: addrspacecast %class.C* %c to %class.C addrspace(4)* + //CHECK: call void @_ZNU3AS41C6AssignE1E(%class.C addrspace(4)* %{{[0-9]+}}, i32 0) + c.Assign(a); + //CHECK: addrspacecast %class.C* %c to %class.C addrspace(4)* + //CHECK: call void @_ZNU3AS41C8OrAssignE1E(%class.C addrspace(4)* %{{[0-9]+}}, i32 0) + c.OrAssign(a); + + E e; + //CHECK: store i32 1, i32* %e + e = b; + //CHECK: store i32 0, i32 addrspace(1)* @globE + globE = a; + //CHECK: store i32 %or, i32 addrspace(1)* @globI + globI |= b; + //CHECK: store i32 %add, i32 addrspace(1)* @globI + globI += a; + //CHECK: store volatile i32 %and, i32 addrspace(1)* @globVI + globVI &= b; + //CHECK: store volatile i32 %sub, i32 addrspace(1)* @globVI + globVI -= a; +} + +//CHECK: define linkonce_odr void @_ZNU3AS41C6AssignE1E(%class.C addrspace(4)* %this, i32 %e) +//CHECK: [[E:%[0-9]+]] = load i32, i32* %e.addr +//CHECK: %me = getelementptr inbounds %class.C, %class.C addrspace(4)* %this1, i32 0, i32 0 +//CHECK: store i32 [[E]], i32 addrspace(4)* %me + +//CHECK define linkonce_odr void @_ZNU3AS41C8OrAssignE1E(%class.C addrspace(4)* %this, i32 %e) +//CHECK: [[E:%[0-9]+]] = load i32, i32* %e.addr +//CHECK: %mi = getelementptr inbounds %class.C, %class.C addrspace(4)* %this1, i32 0, i32 1 +//CHECK: [[MI:%[0-9]+]] = load i32, i32 addrspace(4)* %mi +//CHECK: %or = or i32 [[MI]], [[E]] diff --git a/test/CodeGenOpenCLCXX/addrspace-references.cl b/test/CodeGenOpenCLCXX/addrspace-references.cl new file mode 100644 index 0000000000..b17e701426 --- /dev/null +++ b/test/CodeGenOpenCLCXX/addrspace-references.cl @@ -0,0 +1,14 @@ +//RUN: %clang_cc1 %s -cl-std=c++ -triple spir -emit-llvm -o - | FileCheck %s + +int bar(const unsigned int &i); +// CHECK-LABEL: define spir_func void @_Z3foov() +void foo() { + // The generic addr space reference parameter object will be bound + // to a temporary value allocated in private addr space. We need an + // addrspacecast before passing the value to the function. + // CHECK: [[REF:%.*]] = alloca i32 + // CHECK: store i32 1, i32* [[REF]] + // CHECK: [[REG:%[0-9]+]] = addrspacecast i32* [[REF]] to i32 addrspace(4)* + // CHECK: call spir_func i32 @_Z3barRU3AS4Kj(i32 addrspace(4)* nonnull dereferenceable(4) [[REG]]) + bar(1); +} diff --git a/test/CodeGenOpenCLCXX/local_addrspace_init.cl b/test/CodeGenOpenCLCXX/local_addrspace_init.cl new file mode 100644 index 0000000000..8f78a35343 --- /dev/null +++ b/test/CodeGenOpenCLCXX/local_addrspace_init.cl @@ -0,0 +1,20 @@ +// RUN: %clang_cc1 %s -triple spir -cl-std=c++ -emit-llvm -O0 -o - | FileCheck %s + +// Test that we don't initialize local address space objects. +//CHECK: @_ZZ4testvE1i = internal addrspace(3) global i32 undef +//CHECK: @_ZZ4testvE2ii = internal addrspace(3) global %class.C undef +class C { + int i; +}; + +kernel void test() { + __local int i; + __local C ii; + // FIXME: In OpenCL C we don't accept initializers for local + // address space variables. User defined initialization could + // make sense, but would it mean that all work items need to + // execute it? Potentially disallowing any initialization would + // make things easier and assingments can be used to set specific + // values. This rules should make it consistent with OpenCL C. + //__local C c(); +} diff --git a/test/CodeGenOpenCLCXX/method-overload-address-space.cl b/test/CodeGenOpenCLCXX/method-overload-address-space.cl new file mode 100644 index 0000000000..0864589b05 --- /dev/null +++ b/test/CodeGenOpenCLCXX/method-overload-address-space.cl @@ -0,0 +1,35 @@ +//RUN: %clang_cc1 %s -triple spir-unknown-unknown -cl-std=c++ -emit-llvm -O0 -o - | FileCheck %s + +struct C { + void foo() __local; + void foo() __global; + void foo(); + void bar(); +}; + +__global C c1; + +__kernel void k() { + __local C c2; + C c3; + __global C &c_ref = c1; + __global C *c_ptr; + + // CHECK: call void @_ZNU3AS11C3fooEv(%struct.C addrspace(1)* + c1.foo(); + // CHECK: call void @_ZNU3AS31C3fooEv(%struct.C addrspace(3)* + c2.foo(); + // CHECK: call void @_ZNU3AS41C3fooEv(%struct.C addrspace(4)* + c3.foo(); + // CHECK: call void @_ZNU3AS11C3fooEv(%struct.C addrspace(1)* + c_ptr->foo(); + // CHECK: void @_ZNU3AS11C3fooEv(%struct.C addrspace(1)* + c_ref.foo(); + + // CHECK: call void @_ZNU3AS41C3barEv(%struct.C addrspace(4)* addrspacecast (%struct.C addrspace(1)* @c1 to %struct.C addrspace(4)*)) + c1.bar(); + //FIXME: Doesn't compile yet + //c_ptr->bar(); + // CHECK: call void @_ZNU3AS41C3barEv(%struct.C addrspace(4)* addrspacecast (%struct.C addrspace(1)* @c1 to %struct.C addrspace(4)*)) + c_ref.bar(); +} |