summaryrefslogtreecommitdiff
diff options
context:
space:
mode:
Diffstat
-rw-r--r--clang/include/clang/Sema/Sema.h4
-rw-r--r--clang/lib/AST/Expr.cpp4
-rw-r--r--clang/lib/CodeGen/CGCall.cpp2
-rw-r--r--clang/lib/Sema/SemaDecl.cpp7
-rw-r--r--clang/lib/Sema/SemaDeclCXX.cpp66
-rw-r--r--clang/lib/Sema/SemaInit.cpp11
-rw-r--r--clang/lib/Sema/SemaType.cpp5
-rw-r--r--clang/test/CodeGenOpenCLCXX/addrspace-of-this.cl130
-rw-r--r--clang/test/SemaOpenCLCXX/address-space-templates.cl4
9 files changed, 143 insertions, 90 deletions
diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h
index c9498c19c21..e5b7465820a 100644
--- a/clang/include/clang/Sema/Sema.h
+++ b/clang/include/clang/Sema/Sema.h
@@ -307,6 +307,10 @@ class Sema {
}
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;
diff --git a/clang/lib/AST/Expr.cpp b/clang/lib/AST/Expr.cpp
index 7f6df179d43..7cdd3b2c2a3 100644
--- a/clang/lib/AST/Expr.cpp
+++ b/clang/lib/AST/Expr.cpp
@@ -1677,10 +1677,10 @@ bool CastExpr::CastConsistency() const {
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;
diff --git a/clang/lib/CodeGen/CGCall.cpp b/clang/lib/CodeGen/CGCall.cpp
index 09116d465e3..455a25434ff 100644
--- a/clang/lib/CodeGen/CGCall.cpp
+++ b/clang/lib/CodeGen/CGCall.cpp
@@ -74,7 +74,7 @@ static CanQualType GetThisType(ASTContext &Context, const CXXRecordDecl *RD,
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));
}
diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp
index 0fb4ba60bf8..e853180d76a 100644
--- a/clang/lib/Sema/SemaDecl.cpp
+++ b/clang/lib/Sema/SemaDecl.cpp
@@ -3192,12 +3192,7 @@ bool Sema::MergeFunctionDecl(FunctionDecl *New, NamedDecl *&OldD,
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);
}
diff --git a/clang/lib/Sema/SemaDeclCXX.cpp b/clang/lib/Sema/SemaDeclCXX.cpp
index c2fa87a0790..43b289d8d0d 100644
--- a/clang/lib/Sema/SemaDeclCXX.cpp
+++ b/clang/lib/Sema/SemaDeclCXX.cpp
@@ -6551,8 +6551,11 @@ void Sema::CheckExplicitlyDefaultedSpecialMember(CXXMethodDecl *MD) {
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;
@@ -6560,7 +6563,7 @@ void Sema::CheckExplicitlyDefaultedSpecialMember(CXXMethodDecl *MD) {
}
// A defaulted special member cannot have cv-qualifiers.
- if (Type->getTypeQuals()) {
+ if (Type->getTypeQuals().hasConst() || Type->getTypeQuals().hasVolatile()) {
if (DeleteOnTypeMismatch)
ShouldDeleteForTypeMismatch = true;
else {
@@ -10910,6 +10913,22 @@ void Sema::CheckImplicitSpecialMemberDeclaration(Scope *S, FunctionDecl *FD) {
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:
@@ -10950,9 +10969,7 @@ CXXConstructorDecl *Sema::DeclareImplicitDefaultConstructor(
/* 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.
@@ -11223,9 +11240,7 @@ CXXDestructorDecl *Sema::DeclareImplicitDestructor(CXXRecordDecl *ClassDecl) {
/* 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.
@@ -11799,6 +11814,10 @@ CXXMethodDecl *Sema::DeclareImplicitCopyAssignment(CXXRecordDecl *ClassDecl) {
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,
@@ -11825,10 +11844,7 @@ CXXMethodDecl *Sema::DeclareImplicitCopyAssignment(CXXRecordDecl *ClassDecl) {
/* 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,
@@ -12312,8 +12328,6 @@ void Sema::DefineImplicitMoveAssignment(SourceLocation CurrentLocation,
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()
@@ -12495,6 +12509,10 @@ CXXConstructorDecl *Sema::DeclareImplicitCopyConstructor(
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,
@@ -12523,11 +12541,7 @@ CXXConstructorDecl *Sema::DeclareImplicitCopyConstructor(
/* 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,
@@ -12624,7 +12638,11 @@ CXXConstructorDecl *Sema::DeclareImplicitMoveConstructor(
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,
@@ -12653,11 +12671,7 @@ CXXConstructorDecl *Sema::DeclareImplicitMoveConstructor(
/* 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,
diff --git a/clang/lib/Sema/SemaInit.cpp b/clang/lib/Sema/SemaInit.cpp
index 80cc4617eb5..10c0c6bf33b 100644
--- a/clang/lib/Sema/SemaInit.cpp
+++ b/clang/lib/Sema/SemaInit.cpp
@@ -4669,11 +4669,22 @@ static void TryReferenceInitializationCore(Sema &S,
// 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).
diff --git a/clang/lib/Sema/SemaType.cpp b/clang/lib/Sema/SemaType.cpp
index bf1b9c66a75..b4c075e9c46 100644
--- a/clang/lib/Sema/SemaType.cpp
+++ b/clang/lib/Sema/SemaType.cpp
@@ -4836,11 +4836,8 @@ static TypeSourceInfo *GetFullTypeForDeclarator(TypeProcessingState &state,
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;
}
diff --git a/clang/test/CodeGenOpenCLCXX/addrspace-of-this.cl b/clang/test/CodeGenOpenCLCXX/addrspace-of-this.cl
index af0e3b1a68c..83af3fbcb3b 100644
--- a/clang/test/CodeGenOpenCLCXX/addrspace-of-this.cl
+++ b/clang/test/CodeGenOpenCLCXX/addrspace-of-this.cl
@@ -9,18 +9,21 @@ class C {
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 @@ __kernel void test__global() {
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 @@ __kernel void test__global() {
// 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 @@ __kernel void test__global() {
// 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 @@ TEST(__local)
// 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]])
diff --git a/clang/test/SemaOpenCLCXX/address-space-templates.cl b/clang/test/SemaOpenCLCXX/address-space-templates.cl
index 80762fce8a9..48fbdc7642d 100644
--- a/clang/test/SemaOpenCLCXX/address-space-templates.cl
+++ b/clang/test/SemaOpenCLCXX/address-space-templates.cl
@@ -4,9 +4,7 @@ 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}}
- // 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>
generated by cgit v1.2.3 (git 2.25.1) at 2024-07-10 08:20:39 +0000