Fix PR20495: correct inference of the CUDA target for implicit members

[rsmith]

lib/Sema/SemaCUDA.cpp

120–121

Optional<CUDAFunctionTarget> InferredTarget;

129

For a non-abstract class, you should visit virtual bases as well as direct bases. (For an abstract class, you should skip direct and indirect virtual bases.)

138–142

These are not necessarily correct; the user might have explicitly defaulted a const/volatile/whatever special member function.

156–158

Declaring implicit special members is done lazily, so emitting diagnostics from here will result in erratic behavior. It would be better to defer the diagnostic until the special member is defined in C++98, and to mark the member as deleted in this case in C++11.

lib/Sema/SemaOverload.cpp

5637

Hmm, why do you need this change?

[eliben]

lib/Sema/SemaCUDA.cpp

156–158

If I just mark the member as deleted in C++11, is there away to still emit the descriptive diagnostic? Just deleting the member doesn't provide the user with any hint of what the underlying problem is.

lib/Sema/SemaOverload.cpp

5637

Consider what happens when the target inference method runs. It looks up the ctors of base classes, which gets to this point. Now, at this point we don't yet know what the caller's (the implicit method being defined) target is, so we can not reason whether there's a target mismatch.

AFAIU this added test solves the problem, since we really cannot reject candidates based on target-ness when declaring an implicit member. The implicit member's target-ness will be determined by the inference method, and the collision test (the one with the Diag you commented on) should take care of mismatches.

Does this make sense?

[rsmith]

lib/Sema/SemaCUDA.cpp

138–142

Hmm, I see you're not actually doing this for any defaulted function, just for an implicit one. In that case, this looks fine, but please rename the function to talk about implicit special members not defaulted ones.

156–158

Yes, see Sema::NoteDeletedFunction.

lib/Sema/SemaOverload.cpp

5637

Thanks, that makes sense. This at least deserves an explanatory comment.

[eliben]

Addresses Richard's review comments

[eliben]

Thanks for the review Richard. I believe I've addressed your comments. PTAL

lib/Sema/SemaCUDA.cpp

120–121

Done.

129

Done.

138–142

Renaming done.

156–158

Thanks. I ended up implanting this into ShouldDeleteSpecialMember, with diagnostics; this is called from NoteDeletedFunction, and it also takes care of actually deleting the member in C++11 mode. Let me know if this makes sense.

lib/Sema/SemaOverload.cpp

5637

Comment added, let me know if you'd want to see more details in it.

[rsmith]

This should have more test coverage: at least cover each different kind of special member function, const/non-const arguments, explicitly-defaulted functions, and cases where a special member of a derived class is deleted because the base class's special member would call functions with a __host__ / __device__ mismatch.

lib/Sema/SemaCUDA.cpp

131–141

For a non-abstract class, you'll collect direct vbases twice here (once from the bases list and once from the vbases list). The easiest thing to do would be to change the condition in the first loop to just

if (!B.isVirtual())

172

This should be a note, rather than an error.

lib/Sema/SemaDeclCXX.cpp

5568–5574

This is SMI.ConstArg.

lib/Sema/SemaOverload.cpp

5638

s/we still don't know/we may not yet know/

[eliben]

Addresses Richard's comments.

I made the mismatch detection more robust - now instead of returning "host", when inference fails due to a conflict it returns an "invalid" target, which then causes all lookups to fail. This also provides a more natural place to rerun inference with Diagnose=true, when we try to diagnose why overload resolution failed.

I added more tests, but for now can't test copy/move assignment operators due to http://llvm.org/bugs/show_bug.cgi?id=20886

[eliben]

Thanks for the review. PTAL

lib/Sema/SemaCUDA.cpp

131–141

Done

172

Done

lib/Sema/SemaDeclCXX.cpp

5568–5574

Done

lib/Sema/SemaOverload.cpp

5638

Done

[eliben]

Friendly ping

[eliben]

Friendly ping!

[pcc]

include/clang/Basic/Attr.td

545

This doesn't seem like it should have a spelling.

lib/Sema/SemaCUDA.cpp

101

Is there anything preventing a special member from being manually marked as global? (Apparently not; this parses without errors with a recent clang:)

struct A {
  __attribute__((global)) A() {}
};

lib/Sema/SemaDeclCXX.cpp

10438

It looks like this is already set by the /*isImplicitlyDeclared=*/true parameter passed to CXXConstructorDecl::Create.

10611

Likewise.

[eliben]

Addressing Peter's comments

[eliben]

Thank you for the review, Peter. PTAL

include/clang/Basic/Attr.td

545

Done

lib/Sema/SemaCUDA.cpp

101

Right. I think this should be detected earlier (when a method is constructed) with an appropriate error. I'll change the assertion to a runtime test for now - but am also adding a TODO to have a more descriptive error earlier.

lib/Sema/SemaDeclCXX.cpp

10438

This was cargo-culted from the creation of implicit default ctor and dtor. Fixed now - I've also fixed the source of the carto-culting :) Thanks for noticing this

10611

Fixed

[pcc]

lib/Sema/SemaDeclCXX.cpp

1

I assume this was a mistake.

test/SemaCUDA/implicit-member-target.cu

147

This should derive from A7_with_copy_assign, right?

151

Please update this test once your other patch lands and add the other tests Richard asked for.

[eliben]

PTAL

lib/Sema/SemaDeclCXX.cpp

1

Yes, sorry. Fixing.

test/SemaCUDA/implicit-member-target.cu

147

Done

151

Done.

Note that in the case of the move constructor, the error message is not great. I presume this is because Sema tries to find the copy ctor after the move ctor doesn't match due to target; then since the move ctor is defined, the copy ctor is deleted - hence the message.

[eliben]

Ping!

Since we're mostly discussing tests now, an LGTM will be appreciated. I can fix/add more tests later as needed. This change fixes completely broken functionality in the current frontend, which causes the frontend to fail compiling real code and even crash in some cases.

[pcc]

LGTM with nit.

test/SemaCUDA/implicit-member-target-collision-cxx11.cu

19

Should be ) not viable}} here and elsewhere in this file.

[eliben]

Committed (with nit addressed)

r218624.

Thanks for the review!

include/clang/Basic/Attr.td

 
 def CUDAHost : InheritableAttr {
   let Spellings = [GNU<"host">];
   let Subjects = SubjectList<[Function]>;
   let LangOpts = [CUDA];
   let Documentation = [Undocumented];
 }
 
+def CUDAInvalidTarget : InheritableAttr {
+  let Spellings = [];
+  let Subjects = SubjectList<[Function]>;
+  let LangOpts = [CUDA];
+  let Documentation = [Undocumented];
+}
+
 def CUDALaunchBounds : InheritableAttr {
   let Spellings = [GNU<"launch_bounds">];
   let Args = [IntArgument<"MaxThreads">, DefaultIntArgument<"MinBlocks", 0>];
   let LangOpts = [CUDA];
   let Subjects = SubjectList<[ObjCMethod, FunctionLike], WarnDiag,
                              "ExpectedFunctionOrMethod">;
   // An AST node is created for this attribute, but is not used by other parts
   // of the compiler. However, this node needs to exist in the AST because

include/clang/Basic/DiagnosticSemaKinds.td

     "candidate %select{function|function|constructor|"
     "function |function |constructor |"
     "constructor (the implicit default constructor)|"
     "constructor (the implicit copy constructor)|"
     "constructor (the implicit move constructor)|"
     "function (the implicit copy assignment operator)|"
     "function (the implicit move assignment operator)|"
     "constructor (inherited)}0 not viable: call to "
-    "%select{__device__|__global__|__host__|__host__ __device__}1 function from"
-    " %select{__device__|__global__|__host__|__host__ __device__}2 function">;
+    "%select{__device__|__global__|__host__|__host__ __device__|invalid}1 function from"
+    " %select{__device__|__global__|__host__|__host__ __device__|invalid}2 function">;
+def note_implicit_member_target_infer_collision : Note<
+    "implicit %select{"
+    "default constructor|"
+    "copy constructor|"
+    "move constructor|"
+    "copy assignment operator|"
+    "move assignment operator|"
+    "destructor}0 inferred target collision: call to both "
+    "%select{__device__|__global__|__host__|__host__ __device__}1 and "
+    "%select{__device__|__global__|__host__|__host__ __device__}2 members">;
 
 def note_ambiguous_type_conversion: Note<
     "because of ambiguity in conversion %diff{of $ to $|between types}0,1">;
 def note_ovl_builtin_binary_candidate : Note<
     "built-in candidate %0">;
 def note_ovl_builtin_unary_candidate : Note<
     "built-in candidate %0">;
 def err_ovl_no_viable_function_in_init : Error<

include/clang/Sema/Sema.h

   ExprResult VerifyBitField(SourceLocation FieldLoc, IdentifierInfo *FieldName,
                             QualType FieldTy, bool IsMsStruct,
                             Expr *BitWidth, bool *ZeroWidth = nullptr);
 
   enum CUDAFunctionTarget {
     CFT_Device,
     CFT_Global,
     CFT_Host,
-    CFT_HostDevice
+    CFT_HostDevice,
+    CFT_InvalidTarget
   };
 
   CUDAFunctionTarget IdentifyCUDATarget(const FunctionDecl *D);
 
   bool CheckCUDATarget(CUDAFunctionTarget CallerTarget,
                        CUDAFunctionTarget CalleeTarget);
 
-  bool CheckCUDATarget(const FunctionDecl *Caller, const FunctionDecl *Callee) {
-    return CheckCUDATarget(IdentifyCUDATarget(Caller),
-                           IdentifyCUDATarget(Callee));
-  }
+  bool CheckCUDATarget(const FunctionDecl *Caller, const FunctionDecl *Callee);
+
+  /// Given a implicit special member, infer its CUDA target from the
+  /// calls it needs to make to underlying base/field special members.
+  /// \param ClassDecl the class for which the member is being created.
+  /// \param CSM the kind of special member.
+  /// \param MemberDecl the special member itself.
+  /// \param ConstRHS true if this is a copy operation with a const object on
+  ///        its RHS.
+  /// \param Diagnose true if this call should emit diagnostics.
+  /// \return true if there was an error inferring.
+  /// The result of this call is implicit CUDA target attribute(s) attached to
+  /// the member declaration.
+  bool inferCUDATargetForImplicitSpecialMember(CXXRecordDecl *ClassDecl,
+                                               CXXSpecialMember CSM,
+                                               CXXMethodDecl *MemberDecl,
+                                               bool ConstRHS,
+                                               bool Diagnose);
 
   /// \name Code completion
   //@{
   /// \brief Describes the context in which code completion occurs.
   enum ParserCompletionContext {
     /// \brief Code completion occurs at top-level or namespace context.
     PCC_Namespace,
     /// \brief Code completion occurs within a class, struct, or union.

lib/Sema/SemaCUDA.cpp

 /// \brief This file implements semantic analysis for CUDA constructs.
 ///
 //===----------------------------------------------------------------------===//
 
 #include "clang/Sema/Sema.h"
 #include "clang/AST/ASTContext.h"
 #include "clang/AST/Decl.h"
 #include "clang/Sema/SemaDiagnostic.h"
+#include "llvm/ADT/Optional.h"
+#include "llvm/ADT/SmallVector.h"
 using namespace clang;
 
 ExprResult Sema::ActOnCUDAExecConfigExpr(Scope *S, SourceLocation LLLLoc,
                                          MultiExprArg ExecConfig,
                                          SourceLocation GGGLoc) {
   FunctionDecl *ConfigDecl = Context.getcudaConfigureCallDecl();
   if (!ConfigDecl)
     return ExprError(Diag(LLLLoc, diag::err_undeclared_var_use)
                      << "cudaConfigureCall");
   QualType ConfigQTy = ConfigDecl->getType();
 
   DeclRefExpr *ConfigDR = new (Context)
       DeclRefExpr(ConfigDecl, false, ConfigQTy, VK_LValue, LLLLoc);
   MarkFunctionReferenced(LLLLoc, ConfigDecl);
 
   return ActOnCallExpr(S, ConfigDR, LLLLoc, ExecConfig, GGGLoc, nullptr,
                        /*IsExecConfig=*/true);
 }
 
 /// IdentifyCUDATarget - Determine the CUDA compilation target for this function
 Sema::CUDAFunctionTarget Sema::IdentifyCUDATarget(const FunctionDecl *D) {
-  // Implicitly declared functions (e.g. copy constructors) are
-  // __host__ __device__
-  if (D->isImplicit())
-    return CFT_HostDevice;
+  if (D->hasAttr<CUDAInvalidTargetAttr>())
+    return CFT_InvalidTarget;
 
   if (D->hasAttr<CUDAGlobalAttr>())
     return CFT_Global;
 
   if (D->hasAttr<CUDADeviceAttr>()) {
     if (D->hasAttr<CUDAHostAttr>())
       return CFT_HostDevice;
     return CFT_Device;
   }
 
   return CFT_Host;
 }
 
+bool Sema::CheckCUDATarget(const FunctionDecl *Caller,
+                           const FunctionDecl *Callee) {
+  return CheckCUDATarget(IdentifyCUDATarget(Caller),
+                         IdentifyCUDATarget(Callee));
+}
+
 bool Sema::CheckCUDATarget(CUDAFunctionTarget CallerTarget,
                            CUDAFunctionTarget CalleeTarget) {
+  // If one of the targets is invalid, the check always fails, no matter what
+  // the other target is.
+  if (CallerTarget == CFT_InvalidTarget || CalleeTarget == CFT_InvalidTarget)
+    return true;
+
   // CUDA B.1.1 "The __device__ qualifier declares a function that is...
   // Callable from the device only."
   if (CallerTarget == CFT_Host && CalleeTarget == CFT_Device)
     return true;
 
   // CUDA B.1.2 "The __global__ qualifier declares a function that is...
   // Callable from the host only."
   // CUDA B.1.3 "The __host__ qualifier declares a function that is...
   // Callable from the host only."
   if ((CallerTarget == CFT_Device || CallerTarget == CFT_Global) &&
       (CalleeTarget == CFT_Host || CalleeTarget == CFT_Global))
     return true;
 
   if (CallerTarget == CFT_HostDevice && CalleeTarget != CFT_HostDevice)
     return true;
 
   return false;
 }
 
+/// When an implicitly-declared special member has to invoke more than one
+/// base/field special member, conflicts may occur in the targets of these
+/// members. For example, if one base's member __host__ and another's is
+/// __device__, it's a conflict.
+/// This function figures out if the given targets \param Target1 and
+/// \param Target2 conflict, and if they do not it fills in
+/// \param ResolvedTarget with a target that resolves for both calls.
+/// \return true if there's a conflict, false otherwise.
+static bool
+resolveCalleeCUDATargetConflict(Sema::CUDAFunctionTarget Target1,
+                                Sema::CUDAFunctionTarget Target2,
+                                Sema::CUDAFunctionTarget *ResolvedTarget) {
+  if (Target1 == Sema::CFT_Global && Target2 == Sema::CFT_Global) {
+    // TODO: this shouldn't happen, really. Methods cannot be marked __global__.
+    // Clang should detect this earlier and produce an error. Then this
+    // condition can be changed to an assertion.
+    return true;
+  }
+
+  if (Target1 == Sema::CFT_HostDevice) {
+    *ResolvedTarget = Target2;
+  } else if (Target2 == Sema::CFT_HostDevice) {
+    *ResolvedTarget = Target1;
+  } else if (Target1 != Target2) {
+    return true;
+  } else {
+    *ResolvedTarget = Target1;
+  }
+
+  return false;
+}
+
+bool Sema::inferCUDATargetForImplicitSpecialMember(CXXRecordDecl *ClassDecl,
+                                                   CXXSpecialMember CSM,
+                                                   CXXMethodDecl *MemberDecl,
+                                                   bool ConstRHS,
+                                                   bool Diagnose) {
+  llvm::Optional<CUDAFunctionTarget> InferredTarget;
+
+  // We're going to invoke special member lookup; mark that these special
+  // members are called from this one, and not from its caller.
+  ContextRAII MethodContext(*this, MemberDecl);
+
+  // Look for special members in base classes that should be invoked from here.
+  // Infer the target of this member base on the ones it should call.
+  // Skip direct and indirect virtual bases for abstract classes.
+  llvm::SmallVector<const CXXBaseSpecifier *, 16> Bases;
+  for (const auto &B : ClassDecl->bases()) {
+    if (!B.isVirtual()) {
+      Bases.push_back(&B);
+    }
+  }
+
+  if (!ClassDecl->isAbstract()) {
+    for (const auto &VB : ClassDecl->vbases()) {
+      Bases.push_back(&VB);
+    }
+  }
+
+  for (const auto *B : Bases) {
+    const RecordType *BaseType = B->getType()->getAs<RecordType>();
+    if (!BaseType) {
+      continue;
+    }
+
+    CXXRecordDecl *BaseClassDecl = cast<CXXRecordDecl>(BaseType->getDecl());
+    Sema::SpecialMemberOverloadResult *SMOR =
+        LookupSpecialMember(BaseClassDecl, CSM,
+                            /* ConstArg */ ConstRHS,
+                            /* VolatileArg */ false,
+                            /* RValueThis */ false,
+                            /* ConstThis */ false,
+                            /* VolatileThis */ false);
+
+    if (!SMOR || !SMOR->getMethod()) {
+      continue;
+    }
+
+    CUDAFunctionTarget BaseMethodTarget = IdentifyCUDATarget(SMOR->getMethod());
+    if (!InferredTarget.hasValue()) {
+      InferredTarget = BaseMethodTarget;
+    } else {
+      bool ResolutionError = resolveCalleeCUDATargetConflict(
+          InferredTarget.getValue(), BaseMethodTarget,
+          InferredTarget.getPointer());
+      if (ResolutionError) {
+        if (Diagnose) {
+          Diag(ClassDecl->getLocation(),
+               diag::note_implicit_member_target_infer_collision)
+              << (unsigned)CSM << InferredTarget.getValue() << BaseMethodTarget;
+        }
+        MemberDecl->addAttr(CUDAInvalidTargetAttr::CreateImplicit(Context));
+        return true;
+      }
+    }
+  }
+
+  // Same as for bases, but now for special members of fields.
+  for (const auto *F : ClassDecl->fields()) {
+    if (F->isInvalidDecl()) {
+      continue;
+    }
+
+    const RecordType *FieldType =
+        Context.getBaseElementType(F->getType())->getAs<RecordType>();
+    if (!FieldType) {
+      continue;
+    }
+
+    CXXRecordDecl *FieldRecDecl = cast<CXXRecordDecl>(FieldType->getDecl());
+    Sema::SpecialMemberOverloadResult *SMOR =
+        LookupSpecialMember(FieldRecDecl, CSM,
+                            /* ConstArg */ ConstRHS && !F->isMutable(),
+                            /* VolatileArg */ false,
+                            /* RValueThis */ false,
+                            /* ConstThis */ false,
+                            /* VolatileThis */ false);
+
+    if (!SMOR || !SMOR->getMethod()) {
+      continue;
+    }
+
+    CUDAFunctionTarget FieldMethodTarget =
+        IdentifyCUDATarget(SMOR->getMethod());
+    if (!InferredTarget.hasValue()) {
+      InferredTarget = FieldMethodTarget;
+    } else {
+      bool ResolutionError = resolveCalleeCUDATargetConflict(
+          InferredTarget.getValue(), FieldMethodTarget,
+          InferredTarget.getPointer());
+      if (ResolutionError) {
+        if (Diagnose) {
+          Diag(ClassDecl->getLocation(),
+               diag::note_implicit_member_target_infer_collision)
+              << (unsigned)CSM << InferredTarget.getValue()
+              << FieldMethodTarget;
+        }
+        MemberDecl->addAttr(CUDAInvalidTargetAttr::CreateImplicit(Context));
+        return true;
+      }
+    }
+  }
+
+  if (InferredTarget.hasValue()) {
+    if (InferredTarget.getValue() == CFT_Device) {
+      MemberDecl->addAttr(CUDADeviceAttr::CreateImplicit(Context));
+    } else if (InferredTarget.getValue() == CFT_Host) {
+      MemberDecl->addAttr(CUDAHostAttr::CreateImplicit(Context));
+    } else {
+      MemberDecl->addAttr(CUDADeviceAttr::CreateImplicit(Context));
+      MemberDecl->addAttr(CUDAHostAttr::CreateImplicit(Context));
+    }
+  } else {
+    // If no target was inferred, mark this member as __host__ __device__;
+    // it's the least restrictive option that can be invoked from any target.
+    MemberDecl->addAttr(CUDADeviceAttr::CreateImplicit(Context));
+    MemberDecl->addAttr(CUDAHostAttr::CreateImplicit(Context));
+  }
+
+  return false;
+}

lib/Sema/SemaDeclCXX.cpp

       return true;
 
   // Per DR1611, do not consider virtual bases of constructors of abstract
   // classes, since we are not going to construct them.
   if (!RD->isAbstract() || !SMI.IsConstructor) {
     for (auto &BI : RD->vbases())
       if (SMI.shouldDeleteForBase(&BI))
         return true;
   }
 
   for (auto *FI : RD->fields())
     if (!FI->isInvalidDecl() && !FI->isUnnamedBitfield() &&
         SMI.shouldDeleteForField(FI))
       return true;
 
   if (SMI.shouldDeleteForAllConstMembers())
     return true;
 
+  if (getLangOpts().CUDA) {
+    // We should delete the special member in CUDA mode if target inference
+    // failed.
+    return inferCUDATargetForImplicitSpecialMember(RD, CSM, MD, SMI.ConstArg,
+                                                   Diagnose);
+  }
+
   return false;
 }
 
 /// Perform lookup for a special member of the specified kind, and determine
 /// whether it is trivial. If the triviality can be determined without the
 /// lookup, skip it. This is intended for use when determining whether a
 /// special member of a containing object is trivial, and thus does not ever
 /// perform overload resolution for default constructors.

     StdNamespace = NamespaceDecl::Create(Context, 
                                          Context.getTranslationUnitDecl(),
                                          /*Inline=*/false,
                                          SourceLocation(), SourceLocation(),
                                          &PP.getIdentifierTable().get("std"),
                                          /*PrevDecl=*/nullptr);
     getStdNamespace()->setImplicit(true);
   }
-  
   return getStdNamespace();
 }
 
 bool Sema::isStdInitializerList(QualType Ty, QualType *Element) {
   assert(getLangOpts().CPlusPlus &&
          "Looking for std::initializer_list outside of C++.");
 
   // We're looking for implicit instantiations of

 CXXConstructorDecl *Sema::DeclareImplicitDefaultConstructor(
                                                      CXXRecordDecl *ClassDecl) {
   // C++ [class.ctor]p5:
   //   A default constructor for a class X is a constructor of class X
   //   that can be called without an argument. If there is no
   //   user-declared constructor for class X, a default constructor is
   //   implicitly declared. An implicitly-declared default constructor
   //   is an inline public member of its class.
-  assert(ClassDecl->needsImplicitDefaultConstructor() && 
+  assert(ClassDecl->needsImplicitDefaultConstructor() &&
          "Should not build implicit default constructor!");
 
   DeclaringSpecialMember DSM(*this, ClassDecl, CXXDefaultConstructor);
   if (DSM.isAlreadyBeingDeclared())
     return nullptr;
 
   bool Constexpr = defaultedSpecialMemberIsConstexpr(*this, ClassDecl,
                                                      CXXDefaultConstructor,
                                                      false);
 
   // Create the actual constructor declaration.
   CanQualType ClassType
     = Context.getCanonicalType(Context.getTypeDeclType(ClassDecl));
   SourceLocation ClassLoc = ClassDecl->getLocation();
   DeclarationName Name
     = Context.DeclarationNames.getCXXConstructorName(ClassType);
   DeclarationNameInfo NameInfo(Name, ClassLoc);
   CXXConstructorDecl *DefaultCon = CXXConstructorDecl::Create(
       Context, ClassDecl, ClassLoc, NameInfo, /*Type*/QualType(),
       /*TInfo=*/nullptr, /*isExplicit=*/false, /*isInline=*/true,
       /*isImplicitlyDeclared=*/true, Constexpr);
   DefaultCon->setAccess(AS_public);
   DefaultCon->setDefaulted();
-  DefaultCon->setImplicit();
+
+  if (getLangOpts().CUDA) {
+    inferCUDATargetForImplicitSpecialMember(ClassDecl, CXXDefaultConstructor,
+                                            DefaultCon,
+                                            /* ConstRHS */ false,
+                                            /* 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));
 
   // We don't need to use SpecialMemberIsTrivial here; triviality for default
   // constructors is easy to compute.
   DefaultCon->setTrivial(ClassDecl->hasTrivialDefaultConstructor());

     = Context.DeclarationNames.getCXXDestructorName(ClassType);
   DeclarationNameInfo NameInfo(Name, ClassLoc);
   CXXDestructorDecl *Destructor
       = CXXDestructorDecl::Create(Context, ClassDecl, ClassLoc, NameInfo,
                                   QualType(), nullptr, /*isInline=*/true,
                                   /*isImplicitlyDeclared=*/true);
   Destructor->setAccess(AS_public);
   Destructor->setDefaulted();
-  Destructor->setImplicit();
+
+  if (getLangOpts().CUDA) {
+    inferCUDATargetForImplicitSpecialMember(ClassDecl, CXXDestructor,
+                                            Destructor,
+                                            /* ConstRHS */ false,
+                                            /* 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));
 
   AddOverriddenMethods(ClassDecl, Destructor);
 
   // We don't need to use SpecialMemberIsTrivial here; triviality for

   CXXMethodDecl *CopyAssignment =
       CXXMethodDecl::Create(Context, ClassDecl, ClassLoc, NameInfo, QualType(),
                             /*TInfo=*/nullptr, /*StorageClass=*/SC_None,
                             /*isInline=*/true, Constexpr, SourceLocation());
   CopyAssignment->setAccess(AS_public);
   CopyAssignment->setDefaulted();
   CopyAssignment->setImplicit();
 
+  if (getLangOpts().CUDA) {
+    inferCUDATargetForImplicitSpecialMember(ClassDecl, CXXCopyAssignment,
+                                            CopyAssignment,
+                                            /* ConstRHS */ Const,
+                                            /* Diagnose */ false);
+  }
+
   // Build an exception specification pointing back at this member.
   FunctionProtoType::ExtProtoInfo EPI =
       getImplicitMethodEPI(*this, CopyAssignment);
   CopyAssignment->setType(Context.getFunctionType(RetType, ArgType, EPI));
 
   // Add the parameter to the operator.
   ParmVarDecl *FromParam = ParmVarDecl::Create(Context, CopyAssignment,
                                                ClassLoc, ClassLoc,

   CXXMethodDecl *MoveAssignment =
       CXXMethodDecl::Create(Context, ClassDecl, ClassLoc, NameInfo, QualType(),
                             /*TInfo=*/nullptr, /*StorageClass=*/SC_None,
                             /*isInline=*/true, Constexpr, SourceLocation());
   MoveAssignment->setAccess(AS_public);
   MoveAssignment->setDefaulted();
   MoveAssignment->setImplicit();
 
+  if (getLangOpts().CUDA) {
+    inferCUDATargetForImplicitSpecialMember(ClassDecl, CXXMoveAssignment,
+                                            MoveAssignment,
+                                            /* ConstRHS */ false,
+                                            /* Diagnose */ false);
+  }
+
   // Build an exception specification pointing back at this member.
   FunctionProtoType::ExtProtoInfo EPI =
       getImplicitMethodEPI(*this, MoveAssignment);
   MoveAssignment->setType(Context.getFunctionType(RetType, ArgType, EPI));
 
   // Add the parameter to the operator.
   ParmVarDecl *FromParam = ParmVarDecl::Create(Context, MoveAssignment,
                                                ClassLoc, ClassLoc,

   return ExceptSpec;
 }
 
 CXXConstructorDecl *Sema::DeclareImplicitCopyConstructor(
                                                     CXXRecordDecl *ClassDecl) {
   // C++ [class.copy]p4:
   //   If the class definition does not explicitly declare a copy
   //   constructor, one is declared implicitly.
   assert(ClassDecl->needsImplicitCopyConstructor());
 
   DeclaringSpecialMember DSM(*this, ClassDecl, CXXCopyConstructor);
   if (DSM.isAlreadyBeingDeclared())
     return nullptr;
 
   QualType ClassType = Context.getTypeDeclType(ClassDecl);
   QualType ArgType = ClassType;
   bool Const = ClassDecl->implicitCopyConstructorHasConstParam();

   //   member of its class.
   CXXConstructorDecl *CopyConstructor = CXXConstructorDecl::Create(
       Context, ClassDecl, ClassLoc, NameInfo, QualType(), /*TInfo=*/nullptr,
       /*isExplicit=*/false, /*isInline=*/true, /*isImplicitlyDeclared=*/true,
       Constexpr);
   CopyConstructor->setAccess(AS_public);
   CopyConstructor->setDefaulted();
 
+  if (getLangOpts().CUDA) {
+    inferCUDATargetForImplicitSpecialMember(ClassDecl, CXXCopyConstructor,
+                                            CopyConstructor,
+                                            /* ConstRHS */ Const,
+                                            /* 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));
 
   // Add the parameter to the constructor.
   ParmVarDecl *FromParam = ParmVarDecl::Create(Context, CopyConstructor,

       // If this is a deleted function, add it anyway. This might be conformant
       // with the standard. This might not. I'm not sure. It might not matter.
       // In particular, the problem is that this function never gets called. It
       // might just be ill-formed because this function attempts to refer to
       // a deleted function here.
       if (Constructor)
         ExceptSpec.CalledDecl(F->getLocation(), Constructor);
     }
   }
 
   return ExceptSpec;
 }
 
 CXXConstructorDecl *Sema::DeclareImplicitMoveConstructor(
                                                     CXXRecordDecl *ClassDecl) {
   assert(ClassDecl->needsImplicitMoveConstructor());
 

   //   member of its class.
   CXXConstructorDecl *MoveConstructor = CXXConstructorDecl::Create(
       Context, ClassDecl, ClassLoc, NameInfo, QualType(), /*TInfo=*/nullptr,
       /*isExplicit=*/false, /*isInline=*/true, /*isImplicitlyDeclared=*/true,
       Constexpr);
   MoveConstructor->setAccess(AS_public);
   MoveConstructor->setDefaulted();
 
+  if (getLangOpts().CUDA) {
+    inferCUDATargetForImplicitSpecialMember(ClassDecl, CXXMoveConstructor,
+                                            MoveConstructor,
+                                            /* ConstRHS */ false,
+                                            /* 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));
 
   // Add the parameter to the constructor.
   ParmVarDecl *FromParam = ParmVarDecl::Create(Context, MoveConstructor,

lib/Sema/SemaOverload.cpp

     Candidate.Viable = false;
     Candidate.FailureKind = ovl_fail_too_few_arguments;
     return;
   }
 
   // (CUDA B.1): Check for invalid calls between targets.
   if (getLangOpts().CUDA)
     if (const FunctionDecl *Caller = dyn_cast<FunctionDecl>(CurContext))
-      if (CheckCUDATarget(Caller, Function)) {
+      // Skip the check for callers that are implicit members, because in this
+      // case we may not yet know what the member's target is; the target is
+      // inferred for the member automatically, based on the bases and fields of
+      // the class.
+      if (!Caller->isImplicit() && CheckCUDATarget(Caller, Function)) {
         Candidate.Viable = false;
         Candidate.FailureKind = ovl_fail_bad_target;
         return;
       }
 
   // Determine the implicit conversion sequences for each of the
   // arguments.
   for (unsigned ArgIdx = 0; ArgIdx < Args.size(); ++ArgIdx) {

 
   Sema::CUDAFunctionTarget CallerTarget = S.IdentifyCUDATarget(Caller),
                            CalleeTarget = S.IdentifyCUDATarget(Callee);
 
   std::string FnDesc;
   OverloadCandidateKind FnKind = ClassifyOverloadCandidate(S, Callee, FnDesc);
 
   S.Diag(Callee->getLocation(), diag::note_ovl_candidate_bad_target)
-      << (unsigned) FnKind << CalleeTarget << CallerTarget;
+      << (unsigned)FnKind << CalleeTarget << CallerTarget;
+
+  // This could be an implicit constructor for which we could not infer the
+  // target due to a collsion. Diagnose that case.
+  CXXMethodDecl *Meth = dyn_cast<CXXMethodDecl>(Callee);
+  if (Meth != nullptr && Meth->isImplicit()) {
+    CXXRecordDecl *ParentClass = Meth->getParent();
+    Sema::CXXSpecialMember CSM;
+
+    switch (FnKind) {
+    default:
+      return;
+    case oc_implicit_default_constructor:
+      CSM = Sema::CXXDefaultConstructor;
+      break;
+    case oc_implicit_copy_constructor:
+      CSM = Sema::CXXCopyConstructor;
+      break;
+    case oc_implicit_move_constructor:
+      CSM = Sema::CXXMoveConstructor;
+      break;
+    case oc_implicit_copy_assignment:
+      CSM = Sema::CXXCopyAssignment;
+      break;
+    case oc_implicit_move_assignment:
+      CSM = Sema::CXXMoveAssignment;
+      break;
+    };
+
+    bool ConstRHS = false;
+    if (Meth->getNumParams()) {
+      if (const ReferenceType *RT =
+              Meth->getParamDecl(0)->getType()->getAs<ReferenceType>()) {
+        ConstRHS = RT->getPointeeType().isConstQualified();
+      }
+    }
+
+    S.inferCUDATargetForImplicitSpecialMember(ParentClass, CSM, Meth,
+                                              /* ConstRHS */ ConstRHS,
+                                              /* Diagnose */ true);
+  }
 }
 
 void DiagnoseFailedEnableIfAttr(Sema &S, OverloadCandidate *Cand) {
   FunctionDecl *Callee = Cand->Function;
   EnableIfAttr *Attr = static_cast<EnableIfAttr*>(Cand->DeductionFailure.Data);
 
   S.Diag(Callee->getLocation(),
          diag::note_ovl_candidate_disabled_by_enable_if_attr)

         return false;
     } 
     else if (TargetTypeIsNonStaticMemberFunction)
       return false;
 
     if (FunctionDecl *FunDecl = dyn_cast<FunctionDecl>(Fn)) {
       if (S.getLangOpts().CUDA)
         if (FunctionDecl *Caller = dyn_cast<FunctionDecl>(S.CurContext))
-          if (S.CheckCUDATarget(Caller, FunDecl))
+          if (!Caller->isImplicit() && S.CheckCUDATarget(Caller, FunDecl))
             return false;
 
       // If any candidate has a placeholder return type, trigger its deduction
       // now.
       if (S.getLangOpts().CPlusPlus14 &&
           FunDecl->getReturnType()->isUndeducedType() &&
           S.DeduceReturnType(FunDecl, SourceExpr->getLocStart(), Complain))
         return false;

test/SemaCUDA/implicit-member-target-collision-cxx11.cu

+// RUN: %clang_cc1 -std=gnu++11 -fsyntax-only -verify %s
+
+#include "Inputs/cuda.h"
+
+//------------------------------------------------------------------------------
+// Test 1: collision between two bases
+
+struct A1_with_host_ctor {
+  A1_with_host_ctor() {}
+};
+
+struct B1_with_device_ctor {
+  __device__ B1_with_device_ctor() {}
+};
+
+struct C1_with_collision : A1_with_host_ctor, B1_with_device_ctor {
+};
+
+// expected-note@-3 {{candidate constructor (the implicit default constructor}} not viable
+// expected-note@-4 {{implicit default constructor inferred target collision: call to both __host__ and __device__ members}}
+// expected-note@-5 {{candidate constructor (the implicit copy constructor}} not viable
+// expected-note@-6 {{candidate constructor (the implicit move constructor}} not viable
+
+void hostfoo1() {
+  C1_with_collision c; // expected-error {{no matching constructor}}
+}
+
+//------------------------------------------------------------------------------
+// Test 2: collision between two fields
+
+struct C2_with_collision {
+  A1_with_host_ctor aa;
+  B1_with_device_ctor bb;
+};
+
+// expected-note@-5 {{candidate constructor (the implicit default constructor}} not viable
+// expected-note@-6 {{implicit default constructor inferred target collision: call to both __host__ and __device__ members}}
+// expected-note@-7 {{candidate constructor (the implicit copy constructor}} not viable
+// expected-note@-8 {{candidate constructor (the implicit move constructor}} not viable
+
+void hostfoo2() {
+  C2_with_collision c; // expected-error {{no matching constructor}}
+}
+
+//------------------------------------------------------------------------------
+// Test 3: collision between a field and a base
+
+struct C3_with_collision : A1_with_host_ctor {
+  B1_with_device_ctor bb;
+};
+
+// expected-note@-4 {{candidate constructor (the implicit default constructor}} not viable
+// expected-note@-5 {{implicit default constructor inferred target collision: call to both __host__ and __device__ members}}
+// expected-note@-6 {{candidate constructor (the implicit copy constructor}} not viable
+// expected-note@-7 {{candidate constructor (the implicit move constructor}} not viable
+
+void hostfoo3() {
+  C3_with_collision c; // expected-error {{no matching constructor}}
+}
+
+//------------------------------------------------------------------------------
+// Test 4: collision on resolving a copy ctor
+
+struct A4_with_host_copy_ctor {
+  A4_with_host_copy_ctor() {}
+  A4_with_host_copy_ctor(const A4_with_host_copy_ctor&) {}
+};
+
+struct B4_with_device_copy_ctor {
+  B4_with_device_copy_ctor() {}
+  __device__ B4_with_device_copy_ctor(const B4_with_device_copy_ctor&) {}
+};
+
+struct C4_with_collision : A4_with_host_copy_ctor, B4_with_device_copy_ctor {
+};
+
+// expected-note@-3 {{candidate constructor (the implicit default constructor}} not viable
+// expected-note@-4 {{implicit copy constructor inferred target collision}}
+// expected-note@-5 {{candidate constructor (the implicit copy constructor}} not viable
+
+void hostfoo4() {
+  C4_with_collision c;
+  C4_with_collision c2 = c; // expected-error {{no matching constructor}}
+}
+
+//------------------------------------------------------------------------------
+// Test 5: collision on resolving a move ctor
+
+struct A5_with_host_move_ctor {
+  A5_with_host_move_ctor() {}
+  A5_with_host_move_ctor(A5_with_host_move_ctor&&) {}
+// expected-note@-1 {{copy constructor is implicitly deleted because 'A5_with_host_move_ctor' has a user-declared move constructor}}
+};
+
+struct B5_with_device_move_ctor {
+  B5_with_device_move_ctor() {}
+  __device__ B5_with_device_move_ctor(B5_with_device_move_ctor&&) {}
+};
+
+struct C5_with_collision : A5_with_host_move_ctor, B5_with_device_move_ctor {
+};
+// expected-note@-2 {{deleted}}
+
+void hostfoo5() {
+  C5_with_collision c;
+  // What happens here:
+  // This tries to find the move ctor. Since the move ctor is deleted due to
+  // collision, it then looks for a copy ctor. But copy ctors are implicitly
+  // deleted when move ctors are declared explicitly.
+  C5_with_collision c2(static_cast<C5_with_collision&&>(c)); // expected-error {{call to implicitly-deleted}}
+}

test/SemaCUDA/implicit-member-target-collision.cu

+// RUN: %clang_cc1 -fsyntax-only -verify %s
+
+#include "Inputs/cuda.h"
+
+//------------------------------------------------------------------------------
+// Test 1: collision between two bases
+
+struct A1_with_host_ctor {
+  A1_with_host_ctor() {}
+};
+
+struct B1_with_device_ctor {
+  __device__ B1_with_device_ctor() {}
+};
+
+struct C1_with_collision : A1_with_host_ctor, B1_with_device_ctor {
+};
+
+// expected-note@-3 {{candidate constructor (the implicit default constructor}} not viable
+// expected-note@-4 {{implicit default constructor inferred target collision: call to both __host__ and __device__ members}}
+// expected-note@-5 {{candidate constructor (the implicit copy constructor}} not viable
+
+void hostfoo1() {
+  C1_with_collision c; // expected-error {{no matching constructor}}
+}
+
+//------------------------------------------------------------------------------
+// Test 2: collision between two fields
+
+struct C2_with_collision {
+  A1_with_host_ctor aa;
+  B1_with_device_ctor bb;
+};
+
+// expected-note@-5 {{candidate constructor (the implicit default constructor}} not viable
+// expected-note@-6 {{implicit default constructor inferred target collision: call to both __host__ and __device__ members}}
+// expected-note@-7 {{candidate constructor (the implicit copy constructor}} not viable
+
+void hostfoo2() {
+  C2_with_collision c; // expected-error {{no matching constructor}}
+
+}
+
+//------------------------------------------------------------------------------
+// Test 3: collision between a field and a base
+
+struct C3_with_collision : A1_with_host_ctor {
+  B1_with_device_ctor bb;
+};
+
+// expected-note@-4 {{candidate constructor (the implicit default constructor}} not viable
+// expected-note@-5 {{implicit default constructor inferred target collision: call to both __host__ and __device__ members}}
+// expected-note@-6 {{candidate constructor (the implicit copy constructor}} not viable
+
+void hostfoo3() {
+  C3_with_collision c; // expected-error {{no matching constructor}}
+}

test/SemaCUDA/implicit-member-target.cu

+// RUN: %clang_cc1 -std=gnu++11 -fsyntax-only -verify %s
+
+#include "Inputs/cuda.h"
+
+//------------------------------------------------------------------------------
+// Test 1: infer default ctor to be host.
+
+struct A1_with_host_ctor {
+  A1_with_host_ctor() {}
+};
+
+// The implicit default constructor is inferred to be host because it only needs
+// to invoke a single host constructor (A1_with_host_ctor's). So we'll encounter
+// an error when calling it from a __device__ function, but not from a __host__
+// function.
+struct B1_with_implicit_default_ctor : A1_with_host_ctor {
+};
+
+// expected-note@-3 {{call to __host__ function from __device__}}
+// expected-note@-4 {{candidate constructor (the implicit copy constructor) not viable}}
+// expected-note@-5 {{candidate constructor (the implicit move constructor) not viable}}
+
+void hostfoo() {
+  B1_with_implicit_default_ctor b;
+}
+
+__device__ void devicefoo() {
+  B1_with_implicit_default_ctor b; // expected-error {{no matching constructor}}
+}
+
+//------------------------------------------------------------------------------
+// Test 2: infer default ctor to be device.
+
+struct A2_with_device_ctor {
+  __device__ A2_with_device_ctor() {}
+};
+
+struct B2_with_implicit_default_ctor : A2_with_device_ctor {
+};
+
+// expected-note@-3 {{call to __device__ function from __host__}}
+// expected-note@-4 {{candidate constructor (the implicit copy constructor) not viable}}
+// expected-note@-5 {{candidate constructor (the implicit move constructor) not viable}}
+
+void hostfoo2() {
+  B2_with_implicit_default_ctor b;  // expected-error {{no matching constructor}}
+}
+
+__device__ void devicefoo2() {
+  B2_with_implicit_default_ctor b;
+}
+
+//------------------------------------------------------------------------------
+// Test 3: infer copy ctor
+
+struct A3_with_device_ctors {
+  __host__ A3_with_device_ctors() {}
+  __device__ A3_with_device_ctors(const A3_with_device_ctors&) {}
+};
+
+struct B3_with_implicit_ctors : A3_with_device_ctors {
+};
+
+// expected-note@-3 {{copy constructor of 'B3_with_implicit_ctors' is implicitly deleted}}
+
+void hostfoo3() {
+  B3_with_implicit_ctors b;  // this is OK because the inferred default ctor
+                             // here is __host__
+  B3_with_implicit_ctors b2 = b; // expected-error {{call to implicitly-deleted copy constructor}}
+
+}
+
+//------------------------------------------------------------------------------
+// Test 4: infer default ctor from a field, not a base
+
+struct A4_with_host_ctor {
+  A4_with_host_ctor() {}
+};
+
+struct B4_with_implicit_default_ctor {
+  A4_with_host_ctor field;
+};
+
+// expected-note@-4 {{call to __host__ function from __device__}}
+// expected-note@-5 {{candidate constructor (the implicit copy constructor) not viable}}
+// expected-note@-6 {{candidate constructor (the implicit move constructor) not viable}}
+
+void hostfoo4() {
+  B4_with_implicit_default_ctor b;
+}
+
+__device__ void devicefoo4() {
+  B4_with_implicit_default_ctor b; // expected-error {{no matching constructor}}
+}
+
+//------------------------------------------------------------------------------
+// Test 5: copy ctor with non-const param
+
+struct A5_copy_ctor_constness {
+  __host__ A5_copy_ctor_constness() {}
+  __host__ A5_copy_ctor_constness(A5_copy_ctor_constness&) {}
+};
+
+struct B5_copy_ctor_constness : A5_copy_ctor_constness {
+};
+
+// expected-note@-3 {{candidate constructor (the implicit copy constructor) not viable: call to __host__ function from __device__ function}}
+// expected-note@-4 {{candidate constructor (the implicit default constructor) not viable}}
+
+void hostfoo5(B5_copy_ctor_constness& b_arg) {
+  B5_copy_ctor_constness b = b_arg;
+}
+
+__device__ void devicefoo5(B5_copy_ctor_constness& b_arg) {
+  B5_copy_ctor_constness b = b_arg; // expected-error {{no matching constructor}}
+}
+
+//------------------------------------------------------------------------------
+// Test 6: explicitly defaulted ctor: since they are spelled out, they have
+// a host/device designation explicitly so no inference needs to be done.
+
+struct A6_with_device_ctor {
+  __device__ A6_with_device_ctor() {}
+};
+
+struct B6_with_defaulted_ctor : A6_with_device_ctor {
+  __host__ B6_with_defaulted_ctor() = default;
+};
+
+// expected-note@-3 {{candidate constructor not viable: call to __host__ function from __device__ function}}
+// expected-note@-5 {{candidate constructor (the implicit copy constructor) not viable}}
+// expected-note@-6 {{candidate constructor (the implicit move constructor) not viable}}
+
+__device__ void devicefoo6() {
+  B6_with_defaulted_ctor b; // expected-error {{no matching constructor}}
+}
+
+//------------------------------------------------------------------------------
+// Test 7: copy assignment operator
+
+struct A7_with_copy_assign {
+  A7_with_copy_assign() {}
+  __device__ A7_with_copy_assign& operator=(const A7_with_copy_assign&) {}
+};
+
+struct B7_with_copy_assign : A7_with_copy_assign {
+};
+
+// expected-note@-3 {{copy assignment operator of 'B7_with_copy_assign' is implicitly deleted}}
+
+void hostfoo7() {
+  B7_with_copy_assign b1, b2;
+  b1 = b2; // expected-error {{object of type 'B7_with_copy_assign' cannot be assigned because its copy assignment operator is implicitly deleted}}
+}
+
+//------------------------------------------------------------------------------
+// Test 8: move assignment operator
+
+// definitions for std::move
+namespace std {
+inline namespace foo {
+template <class T> struct remove_reference { typedef T type; };
+template <class T> struct remove_reference<T&> { typedef T type; };
+template <class T> struct remove_reference<T&&> { typedef T type; };
+
+template <class T> typename remove_reference<T>::type&& move(T&& t);
+}
+}
+
+struct A8_with_move_assign {
+  A8_with_move_assign() {}
+  __device__ A8_with_move_assign& operator=(A8_with_move_assign&&) {}
+  __device__ A8_with_move_assign& operator=(const A8_with_move_assign&) {}
+};
+
+struct B8_with_move_assign : A8_with_move_assign {
+};
+
+// expected-note@-3 {{copy assignment operator of 'B8_with_move_assign' is implicitly deleted because base class 'A8_with_move_assign' has no copy assignment operator}}
+
+void hostfoo8() {
+  B8_with_move_assign b1, b2;
+  b1 = std::move(b2); // expected-error {{object of type 'B8_with_move_assign' cannot be assigned because its copy assignment operator is implicitly deleted}}
+}