r352617 - [OpenCL] Add generic addr space to the return of implicit assignment.
Anastasia Stulova via cfe-commits
cfe-commits at lists.llvm.org
Wed Jan 30 03:18:08 PST 2019
Author: stulova
Date: Wed Jan 30 03:18:08 2019
New Revision: 352617
URL: http://llvm.org/viewvc/llvm-project?rev=352617&view=rev
Log:
[OpenCL] Add generic addr space to the return of implicit assignment.
When creating the prototype of implicit assignment operators the
returned reference to the class should be qualified with the same
addr space as 'this' (i.e. __generic in OpenCL).
Differential Revision: https://reviews.llvm.org/D57101
Modified:
cfe/trunk/lib/Sema/SemaDeclCXX.cpp
cfe/trunk/test/CodeGenOpenCLCXX/addrspace-of-this.cl
cfe/trunk/test/SemaOpenCLCXX/address_space_overloading.cl
Modified: cfe/trunk/lib/Sema/SemaDeclCXX.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaDeclCXX.cpp?rev=352617&r1=352616&r2=352617&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaDeclCXX.cpp (original)
+++ cfe/trunk/lib/Sema/SemaDeclCXX.cpp Wed Jan 30 03:18:08 2019
@@ -11813,14 +11813,13 @@ CXXMethodDecl *Sema::DeclareImplicitCopy
return nullptr;
QualType ArgType = Context.getTypeDeclType(ClassDecl);
+ if (Context.getLangOpts().OpenCLCPlusPlus)
+ ArgType = Context.getAddrSpaceQualType(ArgType, LangAS::opencl_generic);
QualType RetType = Context.getLValueReferenceType(ArgType);
bool Const = ClassDecl->implicitCopyAssignmentHasConstParam();
if (Const)
ArgType = ArgType.withConst();
- if (Context.getLangOpts().OpenCLCPlusPlus)
- ArgType = Context.getAddrSpaceQualType(ArgType, LangAS::opencl_generic);
-
ArgType = Context.getLValueReferenceType(ArgType);
bool Constexpr = defaultedSpecialMemberIsConstexpr(*this, ClassDecl,
@@ -12138,6 +12137,8 @@ CXXMethodDecl *Sema::DeclareImplicitMove
// constructor rules.
QualType ArgType = Context.getTypeDeclType(ClassDecl);
+ if (Context.getLangOpts().OpenCLCPlusPlus)
+ ArgType = Context.getAddrSpaceQualType(ArgType, LangAS::opencl_generic);
QualType RetType = Context.getLValueReferenceType(ArgType);
ArgType = Context.getRValueReferenceType(ArgType);
Modified: cfe/trunk/test/CodeGenOpenCLCXX/addrspace-of-this.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCLCXX/addrspace-of-this.cl?rev=352617&r1=352616&r2=352617&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenOpenCLCXX/addrspace-of-this.cl (original)
+++ cfe/trunk/test/CodeGenOpenCLCXX/addrspace-of-this.cl Wed Jan 30 03:18:08 2019
@@ -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;
}
@@ -49,58 +68,76 @@ __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 = 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: %call1 = 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) %call4)
+// 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 +149,62 @@ __kernel void test__global() {
TEST(__local)
-// CHECK-LABEL: _Z11test__localv
-// CHECK: @__cxa_guard_acquire
+// COMMON-LABEL: _Z11test__localv
+// EXPL: @__cxa_guard_acquire
// 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)*))
+// EXPL: 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 = 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 = 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: %call1 = 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]]
Modified: cfe/trunk/test/SemaOpenCLCXX/address_space_overloading.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaOpenCLCXX/address_space_overloading.cl?rev=352617&r1=352616&r2=352617&view=diff
==============================================================================
--- cfe/trunk/test/SemaOpenCLCXX/address_space_overloading.cl (original)
+++ cfe/trunk/test/SemaOpenCLCXX/address_space_overloading.cl Wed Jan 30 03:18:08 2019
@@ -1,12 +1,11 @@
// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only -cl-std=c++
-
-// FIXME: This test shouldn't trigger any errors.
+// expected-no-diagnostics
struct RetGlob {
int dummy;
};
-struct RetGen { //expected-error{{binding value of type '__generic RetGen' to reference to type 'RetGen' drops <<ERROR>> qualifiers}}
+struct RetGen {
char dummy;
};
@@ -19,5 +18,5 @@ void kernel k() {
__local int *ArgLoc;
RetGlob TestGlob = foo(ArgGlob);
RetGen TestGen = foo(ArgGen);
- TestGen = foo(ArgLoc); //expected-note{{in implicit copy assignment operator for 'RetGen' first required here}}
+ TestGen = foo(ArgLoc);
}
More information about the cfe-commits
mailing list