Index: cfe/trunk/include/clang/Sema/Sema.h =================================================================== --- cfe/trunk/include/clang/Sema/Sema.h +++ cfe/trunk/include/clang/Sema/Sema.h @@ -307,6 +307,10 @@ } bool shouldLinkPossiblyHiddenDecl(LookupResult &Old, const NamedDecl *New); + void setupImplicitSpecialMemberType(CXXMethodDecl *SpecialMem, + QualType ResultTy, + ArrayRef Args); + public: typedef OpaquePtr DeclGroupPtrTy; typedef OpaquePtr TemplateTy; Index: cfe/trunk/lib/AST/Expr.cpp =================================================================== --- cfe/trunk/lib/AST/Expr.cpp +++ cfe/trunk/lib/AST/Expr.cpp @@ -1677,10 +1677,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: cfe/trunk/lib/CodeGen/CGCall.cpp =================================================================== --- cfe/trunk/lib/CodeGen/CGCall.cpp +++ cfe/trunk/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: cfe/trunk/lib/Sema/SemaDecl.cpp =================================================================== --- cfe/trunk/lib/Sema/SemaDecl.cpp +++ cfe/trunk/lib/Sema/SemaDecl.cpp @@ -3192,12 +3192,7 @@ if (RequiresAdjustment) { const FunctionType *AdjustedType = New->getType()->getAs(); 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(NewQType); } Index: cfe/trunk/lib/Sema/SemaDeclCXX.cpp =================================================================== --- cfe/trunk/lib/Sema/SemaDeclCXX.cpp +++ cfe/trunk/lib/Sema/SemaDeclCXX.cpp @@ -6551,8 +6551,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; @@ -6560,7 +6563,7 @@ } // 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 @@ CheckFunctionDeclaration(S, FD, R, /*IsMemberSpecialization*/false); } +void Sema::setupImplicitSpecialMemberType(CXXMethodDecl *SpecialMem, + QualType ResultTy, + ArrayRef 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 @@ /* 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 @@ /* 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 @@ 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 @@ /* 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 @@ ParmVarDecl *Other = MoveAssignOperator->getParamDecl(0); QualType OtherRefType = Other->getType()-> getAs()->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 @@ 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 @@ /* 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 @@ 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 @@ /* 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: cfe/trunk/lib/Sema/SemaInit.cpp =================================================================== --- cfe/trunk/lib/Sema/SemaInit.cpp +++ cfe/trunk/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: cfe/trunk/lib/Sema/SemaType.cpp =================================================================== --- cfe/trunk/lib/Sema/SemaType.cpp +++ cfe/trunk/lib/Sema/SemaType.cpp @@ -4836,11 +4836,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: cfe/trunk/test/CodeGenOpenCLCXX/addrspace-of-this.cl =================================================================== --- cfe/trunk/test/CodeGenOpenCLCXX/addrspace-of-this.cl +++ cfe/trunk/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: cfe/trunk/test/SemaOpenCLCXX/address-space-templates.cl =================================================================== --- cfe/trunk/test/SemaOpenCLCXX/address-space-templates.cl +++ cfe/trunk/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