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

[rsmith]

lib/Sema/SemaCUDA.cpp

110–111

Optional<CUDAFunctionTarget> InferredTarget;

119

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.)

128–132

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

146–148

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

146–148

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

128–132

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.

146–148

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

110–111

Done.

119

Done.

128–132

Renaming done.

146–148

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

121–131

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())

162

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

121–131

Done

162

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 ↗

(On Diff #13473)

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

lib/Sema/SemaCUDA.cpp

91

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 ↗

(On Diff #13473)

Done

lib/Sema/SemaCUDA.cpp

91

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 ↗

(On Diff #14090)

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

[eliben]

Committed (with nit addressed)

r218624.

Thanks for the review!

include/clang/Basic/DiagnosticSemaKinds.td

     "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">;
+def err_implicit_member_target_infer_collision : Error<
+    "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

     CFT_HostDevice
   };
 
   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 defaulted (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.
+  /// The result of this call is implicit CUDA target attribute(s) attached to
+  /// the member declaration.
+  void inferCUDATargetForDefaultedSpecialMember(CXXRecordDecl *ClassDecl,
+                                                CXXSpecialMember CSM,
+                                                CXXMethodDecl *MemberDecl,
+                                                bool ConstRHS);
 
   /// \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

ExprResult Sema::ActOnCUDAExecConfigExpr(Scope *S, SourceLocation 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<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) {
   // 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) {
+  assert((Target1 != Sema::CFT_Global && Target2 != Sema::CFT_Global) &&
+         "Special members cannot be marked global");
+
+  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 true;
+}
+
+void Sema::inferCUDATargetForDefaultedSpecialMember(CXXRecordDecl *ClassDecl,
+                                                    CXXSpecialMember CSM,
+                                                    CXXMethodDecl *MemberDecl,
+                                                    bool ConstRHS) {
+  CUDAFunctionTarget InferredTarget;
+  bool HasInferredTarget = false;
+
+  // 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.
+  for (const auto &B : ClassDecl->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 (!HasInferredTarget) {
+      HasInferredTarget = true;
+      InferredTarget = BaseMethodTarget;
+    } else {
+      bool ResolutionError = resolveCalleeCUDATargetConflict(
+          InferredTarget, BaseMethodTarget, &InferredTarget);
+      if (ResolutionError) {
+        Diag(ClassDecl->getLocation(),
+             diag::err_implicit_member_target_infer_collision)
+            << (unsigned)CSM << InferredTarget << BaseMethodTarget;
+        return;
+      }
+    }
+  }
+
+  // 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 (!HasInferredTarget) {
+      HasInferredTarget = true;
+      InferredTarget = FieldMethodTarget;
+    } else {
+      bool ResolutionError = resolveCalleeCUDATargetConflict(
+          InferredTarget, FieldMethodTarget, &InferredTarget);
+      if (ResolutionError) {
+        Diag(ClassDecl->getLocation(),
+             diag::err_implicit_member_target_infer_collision)
+            << (unsigned)CSM << InferredTarget << FieldMethodTarget;
+        return;
+      }
+    }
+  }
+
+  if (HasInferredTarget) {
+    if (InferredTarget == CFT_Device) {
+      MemberDecl->addAttr(CUDADeviceAttr::CreateImplicit(Context));
+    } else if (InferredTarget == 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));
+  }
+}

lib/Sema/SemaDeclCXX.cpp

 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,

   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) {
+    // This has to happen before ShouldDeleteSpecialMember is called. In the
+    // absence of this inference, ShouldDeleteSpecialMember may wrongly decide
+    // to delete this ctor.
+    inferCUDATargetForDefaultedSpecialMember(ClassDecl, CXXDefaultConstructor,
+                                             DefaultCon,
+                                             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());
 

   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) {
+    inferCUDATargetForDefaultedSpecialMember(ClassDecl, CXXDestructor,
+                                             Destructor,
+                                             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
   // destructors is easy to compute.

   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) {
+    inferCUDATargetForDefaultedSpecialMember(ClassDecl, CXXCopyAssignment,
+                                             CopyAssignment,
+                                             Const);
+  }
+
   // 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) {
+    inferCUDATargetForDefaultedSpecialMember(ClassDecl, CXXMoveAssignment,
+                                             MoveAssignment,
+                                             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,

   //   An implicitly-declared copy constructor is an inline public
   //   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();
+  CopyConstructor->setImplicit();
+
+  if (getLangOpts().CUDA) {
+    inferCUDATargetForDefaultedSpecialMember(ClassDecl, CXXCopyConstructor,
+                                             CopyConstructor,
+                                             Const);
+  }
 
   // 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,
                                                ClassLoc, ClassLoc,
                                                /*IdentifierInfo=*/nullptr,
                                                ArgType, /*TInfo=*/nullptr,
                                                SC_None, nullptr);

   //   An implicitly-declared copy/move constructor is an inline public
   //   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();
+  MoveConstructor->setImplicit();
+
+  if (getLangOpts().CUDA) {
+    inferCUDATargetForDefaultedSpecialMember(ClassDecl, CXXMoveConstructor,
+                                             MoveConstructor,
+                                             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,
                                                ClassLoc, ClassLoc,
                                                /*IdentifierInfo=*/nullptr,
                                                ArgType, /*TInfo=*/nullptr,
                                                SC_None, nullptr);
   MoveConstructor->setParams(FromParam);

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)) {
+      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) {
     if (ArgIdx < NumParams) {

         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.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-error@-3 {{implicit default constructor inferred target collision: call to both __host__ and __device__ members}}
+// expected-error@-4 {{implicit copy constructor inferred target collision: call to both __host__ __device__ and __host__ __device__ members}}
+
+void hostfoo1() {
+  C1_with_collision c;
+}
+
+//------------------------------------------------------------------------------
+// Test 2: collision between two fields
+
+struct C2_with_collision {
+  A1_with_host_ctor aa;
+  B1_with_device_ctor bb;
+};
+
+// expected-error@-5 {{implicit default constructor inferred target collision: call to both __host__ and __device__ members}}
+// expected-error@-6 {{implicit copy constructor inferred target collision: call to both __host__ __device__ and __host__ __device__ members}}
+
+void hostfoo2() {
+  C2_with_collision c;
+}
+
+//------------------------------------------------------------------------------
+// Test 3: collision between a field and a base
+
+struct C3_with_collision : A1_with_host_ctor {
+  B1_with_device_ctor bb;
+};
+
+// expected-error@-4 {{implicit default constructor inferred target collision: call to both __host__ and __device__ members}}
+// expected-error@-5 {{implicit copy constructor inferred target collision: call to both __host__ __device__ and __host__ __device__ members}}
+
+void hostfoo4() {
+  C3_with_collision c;
+}

test/SemaCUDA/implicit-member-target.cu

+// RUN: %clang_cc1 -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 {{requires 1 argument}}
+
+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 {{requires 1 argument}}
+
+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 {{(the implicit copy constructor) not viable: call to __device__ function from __host__}}
+// expected-note@-4 {{requires 0 arguments}}
+
+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 {{no matching 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 {{requires 1 argument}}
+
+void hostfoo4() {
+  B4_with_implicit_default_ctor b;
+}
+
+__device__ void devicefoo4() {
+  B4_with_implicit_default_ctor b; // expected-error {{no matching constructor}}
+}