Anastasia updated this revision to Diff 181256.
Anastasia added a comment.
- Set correctly value kind of address space conversion after materializing the
temporary
- Removed changes in various places including CodeGen for address space
conversion that became unnecessary after the AST is generated correctly.
- Updated tests to check for the right IR and also not to use hard-coded
registers.
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D56066/new/
https://reviews.llvm.org/D56066
Files:
include/clang/Sema/Sema.h
lib/AST/Expr.cpp
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
@@ -4,9 +4,7 @@
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}}
- // 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()
@@ -69,17 +70,36 @@
// CHECK: %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: %0 = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
-// CHECK: call void @_ZNU3AS41CC1ERU3AS4KS_(%class.C addrspace(4)* %0, %class.C addrspace(4)* dereferenceable(4) addrspacecast (%class.C addrspace(1)* @c to %class.C addrspace(4)*))
+// 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)*))
// 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: [[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 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)
+// 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]])
+
+// 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]])
+
+// 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)
+
+// 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)
+
#define TEST(AS) \
__kernel void test##AS() { \
@@ -95,60 +115,74 @@
// CHECK-LABEL: _Z11test__localv
// CHECK: @__cxa_guard_acquire
-// Test the address space of 'this' when invoking a method.
+// 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 the address space of 'this' when invoking copy-constructor.
+// 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)*))
+
+// 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)*))
+
// Test the address space of 'this' when invoking a constructor.
-// CHECK: %3 = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
-// CHECK: call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* %3)
+// 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 assignment operator.
-// CHECK: %4 = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
-// CHECK: %5 = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
-// CHECK: %call1 = call dereferenceable(4) %class.C addrspace(4)* @_ZNU3AS41CaSERU3AS4KS_(%class.C addrspace(4)* %5, %class.C addrspace(4)* dereferenceable(4) %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: %call1 = call dereferenceable(4) %class.C addrspace(4)* @_ZNU3AS41CaSERU3AS4KS_(%class.C addrspace(4)* [[C2GEN]], %class.C addrspace(4)* dereferenceable(4) [[C1GEN]])
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: %1 = addrspacecast %class.C* %c to %class.C addrspace(4)*
-// CHECK: %call = call i32 @_ZNU3AS41C3getEv(%class.C addrspace(4)* %1)
+// 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: %2 = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
-// CHECK: %3 = addrspacecast %class.C* %c to %class.C addrspace(4)*
-// CHECK: call void @_ZNU3AS41CC1ERU3AS4KS_(%class.C addrspace(4)* %2, %class.C addrspace(4)* dereferenceable(4) %3)
+// 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: %4 = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
-// CHECK: call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* %4)
+// 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: %5 = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
-// CHECK: %6 = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
-// CHECK: %call1 = call dereferenceable(4) %class.C addrspace(4)* @_ZNU3AS41CaSERU3AS4KS_(%class.C addrspace(4)* %6, %class.C addrspace(4)* dereferenceable(4) %5)
+// 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]])
+
// 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: [[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: %2 = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
-// CHECK: %3 = addrspacecast %class.C* %c to %class.C addrspace(4)*
-// CHECK: call void @_ZNU3AS41CC1ERU3AS4KS_(%class.C addrspace(4)* %2, %class.C addrspace(4)* dereferenceable(4) %3)
+// 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: %4 = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
-// CHECK: call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* %4)
+// 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: %5 = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
-// CHECK: %6 = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
-// CHECK: %call1 = call dereferenceable(4) %class.C addrspace(4)* @_ZNU3AS41CaSERU3AS4KS_(%class.C addrspace(4)* %6, %class.C addrspace(4)* dereferenceable(4) %5)
+// 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]])
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
@@ -4669,11 +4669,22 @@
// If the converted initializer is a prvalue, its type T4 is adjusted
// to type "cv1 T4" and the temporary materialization conversion is
// applied.
+ // Postpone address space conversions to after the temporary materialization
+ // conversion to allow creating temporaries in the alloca address space.
+ auto AS1 = T1Quals.getAddressSpace();
+ auto AS2 = T2Quals.getAddressSpace();
+ T1Quals.removeAddressSpace();
+ T2Quals.removeAddressSpace();
QualType cv1T4 = S.Context.getQualifiedType(cv2T2, T1Quals);
if (T1Quals != T2Quals)
Sequence.AddQualificationConversionStep(cv1T4, ValueKind);
Sequence.AddReferenceBindingStep(cv1T4, ValueKind == VK_RValue);
ValueKind = isLValueRef ? VK_LValue : VK_XValue;
+ if (AS1 != AS2) {
+ T1Quals.addAddressSpace(AS1);
+ QualType cv1AST4 = S.Context.getQualifiedType(cv2T2, T1Quals);
+ Sequence.AddQualificationConversionStep(cv1AST4, ValueKind);
+ }
// In any case, the reference is bound to the resulting glvalue (or to
// an appropriate base class subobject).
Index: lib/Sema/SemaDeclCXX.cpp
===================================================================
--- lib/Sema/SemaDeclCXX.cpp
+++ lib/Sema/SemaDeclCXX.cpp
@@ -6546,8 +6546,11 @@
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 +6558,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 +10916,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 +10972,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 +11243,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 +11817,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 +11847,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 +12331,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 +12512,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 +12544,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 +12641,11 @@
return nullptr;
QualType ClassType = Context.getTypeDeclType(ClassDecl);
- QualType ArgType = Context.getRValueReferenceType(ClassType);
+
+ QualType ArgType = ClassType;
+ if (Context.getLangOpts().OpenCLCPlusPlus)
+ ArgType = Context.getAddrSpaceQualType(ClassType, LangAS::opencl_generic);
+ ArgType = Context.getRValueReferenceType(ArgType);
bool Constexpr = defaultedSpecialMemberIsConstexpr(*this, ClassDecl,
CXXMoveConstructor,
@@ -12656,11 +12674,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,7 @@
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: lib/AST/Expr.cpp
===================================================================
--- lib/AST/Expr.cpp
+++ lib/AST/Expr.cpp
@@ -1615,10 +1615,10 @@
auto Ty = getType();
auto SETy = getSubExpr()->getType();
assert(getValueKindForType(Ty) == Expr::getValueKindForType(SETy));
- if (!isGLValue())
+ if (isRValue()) {
Ty = Ty->getPointeeType();
- if (!isGLValue())
SETy = SETy->getPointeeType();
+ }
assert(!Ty.isNull() && !SETy.isNull() &&
Ty.getAddressSpace() != SETy.getAddressSpace());
goto CheckNoBasePath;
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
