Anastasia created this revision.
Anastasia added reviewers: rjmccall, ebevhan.
Herald added a subscriber: yaxunl.
Fixed several issues reported in https://reviews.llvm.org/D55656 and
https://reviews.llvm.org/D54862 - mainly removed addr space from the `QualType`
of the method itself.
This should make all implicitly and explicity defaulted special class members
work now!
Since Mikael is on holiday and I have no permission to upload to his review, I
have to create a new one.
https://reviews.llvm.org/D56066
Files:
include/clang/Sema/Sema.h
lib/CodeGen/CGCall.cpp
lib/Sema/SemaDecl.cpp
lib/Sema/SemaDeclCXX.cpp
lib/Sema/SemaInit.cpp
lib/Sema/SemaType.cpp
test/CodeGenOpenCLCXX/addrspace-of-this.cl
test/SemaOpenCLCXX/address-space-templates.cl
Index: test/SemaOpenCLCXX/address-space-templates.cl
===================================================================
--- test/SemaOpenCLCXX/address-space-templates.cl
+++ test/SemaOpenCLCXX/address-space-templates.cl
@@ -2,11 +2,10 @@
template <typename T>
struct S {
- T a; // expected-error{{field may not be qualified with an address space}}
- T f1(); // expected-error{{function type may not be qualified with an address space}}
+ T a; // expected-error{{field may not be qualified with an address space}}
+ T f1(); // expected-error{{function type may not be qualified with an address space}}
// FIXME: Should only get the error message once.
- void f2(T); // expected-error{{parameter may not be qualified with an address space}} expected-error{{parameter may not be qualified with an address space}}
-
+ void f2(T); // expected-error{{parameter may not be qualified with an address space}}
};
template <typename T>
Index: test/CodeGenOpenCLCXX/addrspace-of-this.cl
===================================================================
--- test/CodeGenOpenCLCXX/addrspace-of-this.cl
+++ test/CodeGenOpenCLCXX/addrspace-of-this.cl
@@ -9,18 +9,21 @@
public:
int v;
C() { v = 2; }
- // FIXME: Does not work yet.
- // C(C &&c) { v = c.v; }
+ C(C &&c) { v = c.v; }
C(const C &c) { v = c.v; }
C &operator=(const C &c) {
v = c.v;
return *this;
}
- // FIXME: Does not work yet.
- //C &operator=(C&& c) & {
- // v = c.v;
- // return *this;
- //}
+ C &operator=(C &&c) & {
+ v = c.v;
+ return *this;
+ }
+
+ C operator+(const C &c) {
+ v += c.v;
+ return *this;
+ }
int get() { return v; }
@@ -41,15 +44,13 @@
C c1(c);
C c2;
c2 = c1;
- // FIXME: Does not work yet.
- // C c3 = c1 + c2;
- // C c4(foo());
- // C c5 = foo();
-
+ C c3 = c1 + c2;
+ C c4(foo());
+ 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)*)) #4
+// 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 the constructor
// CHECK-LABEL: @_ZNU3AS41CC1Ev(%class.C addrspace(4)* %this)
@@ -57,7 +58,7 @@
// 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) #4
+// CHECK: call void @_ZNU3AS41CC2Ev(%class.C addrspace(4)* %this1)
// CHECK: ret void
// CHECK-LABEL: @_Z12test__globalv()
@@ -74,13 +75,28 @@
// Test the address space of 'this' when invoking a constructor.
// CHECK: %1 = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
-// CHECK: call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* %1) #4
+// CHECK: call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* %1)
// Test the address space of 'this' when invoking assignment operator.
// CHECK: %2 = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
// CHECK: %3 = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
// CHECK: %call2 = call dereferenceable(4) %class.C addrspace(4)* @_ZNU3AS41CaSERU3AS4KS_(%class.C addrspace(4)* %3, %class.C addrspace(4)* dereferenceable(4) %2)
+// Test the address space of 'this' when invoking the operator+
+// CHECK: %4 = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
+// CHECK: %5 = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
+// CHECK: call void @_ZNU3AS41CplERU3AS4KS_(%class.C* sret %c3, %class.C addrspace(4)* %4, %class.C addrspace(4)* dereferenceable(4) %5)
+
+// Test the address space of 'this' when invoking the move constructor
+// CHECK: %6 = 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)* %6, %class.C addrspace(4)* dereferenceable(4) %call3)
+
+// Test the address space of 'this' when invoking the move assignment
+// CHECK: %7 = 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)* %7, %class.C addrspace(4)* dereferenceable(4) %call4)
+
#define TEST(AS) \
__kernel void test##AS() { \
AS C c; \
@@ -137,7 +153,7 @@
// CHECK-LABEL: @_Z4testv()
// Test the address space of 'this' when invoking a method.
// CHECK: %1 = addrspacecast %class.C* %c to %class.C addrspace(4)*
-// CHECK: %call = call i32 @_ZNU3AS41C3getEv(%class.C addrspace(4)* %1) #4
+// CHECK: %call = call i32 @_ZNU3AS41C3getEv(%class.C addrspace(4)* %1)
// Test the address space of 'this' when invoking a copy-constructor.
// CHECK: %2 = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
Index: lib/Sema/SemaType.cpp
===================================================================
--- lib/Sema/SemaType.cpp
+++ lib/Sema/SemaType.cpp
@@ -4839,11 +4839,8 @@
LangAS AS =
(CurAS == LangAS::Default ? LangAS::opencl_generic : CurAS);
EPI.TypeQuals.addAddressSpace(AS);
- T = Context.getFunctionType(T, ParamTys, EPI);
- T = state.getSema().Context.getAddrSpaceQualType(T, AS);
- } else {
- T = Context.getFunctionType(T, ParamTys, EPI);
}
+ T = Context.getFunctionType(T, ParamTys, EPI);
}
break;
}
Index: lib/Sema/SemaInit.cpp
===================================================================
--- lib/Sema/SemaInit.cpp
+++ lib/Sema/SemaInit.cpp
@@ -4535,6 +4535,9 @@
= S.CompareReferenceRelationship(DeclLoc, cv1T1, cv2T2, DerivedToBase,
ObjCConversion, ObjCLifetimeConversion);
+ if (InitCategory.isPRValue() || InitCategory.isXValue())
+ T1Quals.removeAddressSpace();
+
// C++0x [dcl.init.ref]p5:
// A reference to type "cv1 T1" is initialized by an expression of type
// "cv2 T2" as follows:
Index: lib/Sema/SemaDeclCXX.cpp
===================================================================
--- lib/Sema/SemaDeclCXX.cpp
+++ lib/Sema/SemaDeclCXX.cpp
@@ -6546,8 +6546,12 @@
if (CSM == CXXCopyAssignment || CSM == CXXMoveAssignment) {
// Check for return type matching.
ReturnType = Type->getReturnType();
- QualType ExpectedReturnType =
- Context.getLValueReferenceType(Context.getTypeDeclType(RD));
+
+ QualType DeclType = Context.getTypeDeclType(RD);
+ DeclType = Context.getAddrSpaceQualType(
+ DeclType, MD->getTypeQualifiers().getAddressSpace());
+ QualType ExpectedReturnType = Context.getLValueReferenceType(DeclType);
+
if (!Context.hasSameType(ReturnType, ExpectedReturnType)) {
Diag(MD->getLocation(), diag::err_defaulted_special_member_return_type)
<< (CSM == CXXMoveAssignment) << ExpectedReturnType;
@@ -6555,7 +6559,7 @@
}
// A defaulted special member cannot have cv-qualifiers.
- if (Type->getTypeQuals()) {
+ if (Type->getTypeQuals().hasConst() || Type->getTypeQuals().hasVolatile()) {
if (DeleteOnTypeMismatch)
ShouldDeleteForTypeMismatch = true;
else {
@@ -10913,6 +10917,22 @@
CheckFunctionDeclaration(S, FD, R, /*IsMemberSpecialization*/false);
}
+void Sema::setupImplicitSpecialMemberType(CXXMethodDecl *SpecialMem,
+ QualType ResultTy,
+ ArrayRef<QualType> Args) {
+ // Build an exception specification pointing back at this constructor.
+ FunctionProtoType::ExtProtoInfo EPI = getImplicitMethodEPI(*this, SpecialMem);
+
+ if (getLangOpts().OpenCLCPlusPlus) {
+ // OpenCL: Implicitly defaulted special member are of the generic address
+ // space.
+ EPI.TypeQuals.addAddressSpace(LangAS::opencl_generic);
+ }
+
+ auto QT = Context.getFunctionType(ResultTy, Args, EPI);
+ SpecialMem->setType(QT);
+}
+
CXXConstructorDecl *Sema::DeclareImplicitDefaultConstructor(
CXXRecordDecl *ClassDecl) {
// C++ [class.ctor]p5:
@@ -10953,9 +10973,7 @@
/* Diagnose */ false);
}
- // Build an exception specification pointing back at this constructor.
- FunctionProtoType::ExtProtoInfo EPI = getImplicitMethodEPI(*this, DefaultCon);
- DefaultCon->setType(Context.getFunctionType(Context.VoidTy, None, EPI));
+ setupImplicitSpecialMemberType(DefaultCon, Context.VoidTy, None);
// We don't need to use SpecialMemberIsTrivial here; triviality for default
// constructors is easy to compute.
@@ -11226,9 +11244,7 @@
/* Diagnose */ false);
}
- // Build an exception specification pointing back at this destructor.
- FunctionProtoType::ExtProtoInfo EPI = getImplicitMethodEPI(*this, Destructor);
- Destructor->setType(Context.getFunctionType(Context.VoidTy, None, EPI));
+ setupImplicitSpecialMemberType(Destructor, Context.VoidTy, None);
// We don't need to use SpecialMemberIsTrivial here; triviality for
// destructors is easy to compute.
@@ -11802,6 +11818,10 @@
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,
@@ -11828,10 +11848,7 @@
/* Diagnose */ false);
}
- // Build an exception specification pointing back at this member.
- FunctionProtoType::ExtProtoInfo EPI =
- getImplicitMethodEPI(*this, CopyAssignment);
- CopyAssignment->setType(Context.getFunctionType(RetType, ArgType, EPI));
+ setupImplicitSpecialMemberType(CopyAssignment, RetType, ArgType);
// Add the parameter to the operator.
ParmVarDecl *FromParam = ParmVarDecl::Create(Context, CopyAssignment,
@@ -12315,8 +12332,6 @@
ParmVarDecl *Other = MoveAssignOperator->getParamDecl(0);
QualType OtherRefType = Other->getType()->
getAs<RValueReferenceType>()->getPointeeType();
- assert(!OtherRefType.getQualifiers() &&
- "Bad argument type of defaulted move assignment");
// Our location for everything implicitly-generated.
SourceLocation Loc = MoveAssignOperator->getEndLoc().isValid()
@@ -12498,6 +12513,10 @@
bool Const = ClassDecl->implicitCopyConstructorHasConstParam();
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,
@@ -12526,11 +12545,7 @@
/* Diagnose */ false);
}
- // Build an exception specification pointing back at this member.
- FunctionProtoType::ExtProtoInfo EPI =
- getImplicitMethodEPI(*this, CopyConstructor);
- CopyConstructor->setType(
- Context.getFunctionType(Context.VoidTy, ArgType, EPI));
+ setupImplicitSpecialMemberType(CopyConstructor, Context.VoidTy, ArgType);
// Add the parameter to the constructor.
ParmVarDecl *FromParam = ParmVarDecl::Create(Context, CopyConstructor,
@@ -12627,7 +12642,11 @@
return nullptr;
QualType ClassType = Context.getTypeDeclType(ClassDecl);
- QualType ArgType = Context.getRValueReferenceType(ClassType);
+
+ QualType ArgType;
+ if (Context.getLangOpts().OpenCLCPlusPlus)
+ ArgType = Context.getAddrSpaceQualType(ClassType, LangAS::opencl_generic);
+ ArgType = Context.getRValueReferenceType(ClassType);
bool Constexpr = defaultedSpecialMemberIsConstexpr(*this, ClassDecl,
CXXMoveConstructor,
@@ -12656,11 +12675,7 @@
/* Diagnose */ false);
}
- // Build an exception specification pointing back at this member.
- FunctionProtoType::ExtProtoInfo EPI =
- getImplicitMethodEPI(*this, MoveConstructor);
- MoveConstructor->setType(
- Context.getFunctionType(Context.VoidTy, ArgType, EPI));
+ setupImplicitSpecialMemberType(MoveConstructor, Context.VoidTy, ArgType);
// Add the parameter to the constructor.
ParmVarDecl *FromParam = ParmVarDecl::Create(Context, MoveConstructor,
Index: lib/Sema/SemaDecl.cpp
===================================================================
--- lib/Sema/SemaDecl.cpp
+++ lib/Sema/SemaDecl.cpp
@@ -3192,12 +3192,7 @@
if (RequiresAdjustment) {
const FunctionType *AdjustedType = New->getType()->getAs<FunctionType>();
AdjustedType = Context.adjustFunctionType(AdjustedType, NewTypeInfo);
-
- QualType AdjustedQT = QualType(AdjustedType, 0);
- LangAS AS = Old->getType().getAddressSpace();
- AdjustedQT = Context.getAddrSpaceQualType(AdjustedQT, AS);
-
- New->setType(AdjustedQT);
+ New->setType(QualType(AdjustedType, 0));
NewQType = Context.getCanonicalType(New->getType());
NewType = cast<FunctionType>(NewQType);
}
Index: lib/CodeGen/CGCall.cpp
===================================================================
--- lib/CodeGen/CGCall.cpp
+++ lib/CodeGen/CGCall.cpp
@@ -74,7 +74,8 @@
const CXXMethodDecl *MD) {
QualType RecTy = Context.getTagDeclType(RD)->getCanonicalTypeInternal();
if (MD)
- RecTy = Context.getAddrSpaceQualType(RecTy, MD->getType().getAddressSpace());
+ RecTy = Context.getAddrSpaceQualType(
+ RecTy, MD->getTypeQualifiers().getAddressSpace());
return Context.getPointerType(CanQualType::CreateUnsafe(RecTy));
}
Index: include/clang/Sema/Sema.h
===================================================================
--- include/clang/Sema/Sema.h
+++ include/clang/Sema/Sema.h
@@ -307,6 +307,10 @@
}
bool shouldLinkPossiblyHiddenDecl(LookupResult &Old, const NamedDecl *New);
+ void setupImplicitSpecialMemberType(CXXMethodDecl *SpecialMem,
+ QualType ResultTy,
+ ArrayRef<QualType> Args);
+
public:
typedef OpaquePtr<DeclGroupRef> DeclGroupPtrTy;
typedef OpaquePtr<TemplateName> TemplateTy;
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
