Index: cfe/trunk/include/clang/Sema/Sema.h
===================================================================
--- cfe/trunk/include/clang/Sema/Sema.h
+++ cfe/trunk/include/clang/Sema/Sema.h
@@ -8325,9 +8325,6 @@
 
   CUDAFunctionTarget IdentifyCUDATarget(const FunctionDecl *D);
 
-  bool CheckCUDATarget(CUDAFunctionTarget CallerTarget,
-                       CUDAFunctionTarget CalleeTarget);
-
   bool CheckCUDATarget(const FunctionDecl *Caller, const FunctionDecl *Callee);
 
   /// Given a implicit special member, infer its CUDA target from the
Index: cfe/trunk/lib/Sema/SemaCUDA.cpp
===================================================================
--- cfe/trunk/lib/Sema/SemaCUDA.cpp
+++ cfe/trunk/lib/Sema/SemaCUDA.cpp
@@ -62,12 +62,9 @@
 
 bool Sema::CheckCUDATarget(const FunctionDecl *Caller,
                            const FunctionDecl *Callee) {
-  return CheckCUDATarget(IdentifyCUDATarget(Caller),
-                         IdentifyCUDATarget(Callee));
-}
+  CUDAFunctionTarget CallerTarget = IdentifyCUDATarget(Caller),
+                     CalleeTarget = 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)
@@ -90,8 +87,11 @@
   // however, in which case the function is compiled for both the host and the
   // device. The __CUDA_ARCH__ macro [...] can be used to differentiate code
   // paths between host and device."
-  bool InDeviceMode = getLangOpts().CUDAIsDevice;
   if (CallerTarget == CFT_HostDevice && CalleeTarget != CFT_HostDevice) {
+    // If the caller is implicit then the check always passes.
+    if (Caller->isImplicit()) return false;
+
+    bool InDeviceMode = getLangOpts().CUDAIsDevice;
     if ((InDeviceMode && CalleeTarget != CFT_Device) ||
         (!InDeviceMode && CalleeTarget != CFT_Host))
       return true;
Index: cfe/trunk/lib/Sema/SemaExpr.cpp
===================================================================
--- cfe/trunk/lib/Sema/SemaExpr.cpp
+++ cfe/trunk/lib/Sema/SemaExpr.cpp
@@ -1590,11 +1590,10 @@
   if (getLangOpts().CUDA)
     if (const FunctionDecl *Caller = dyn_cast<FunctionDecl>(CurContext))
       if (const FunctionDecl *Callee = dyn_cast<FunctionDecl>(D)) {
-        CUDAFunctionTarget CallerTarget = IdentifyCUDATarget(Caller),
-                           CalleeTarget = IdentifyCUDATarget(Callee);
-        if (CheckCUDATarget(CallerTarget, CalleeTarget)) {
+        if (CheckCUDATarget(Caller, Callee)) {
           Diag(NameInfo.getLoc(), diag::err_ref_bad_target)
-            << CalleeTarget << D->getIdentifier() << CallerTarget;
+            << IdentifyCUDATarget(Callee) << D->getIdentifier()
+            << IdentifyCUDATarget(Caller);
           Diag(D->getLocation(), diag::note_previous_decl)
             << D->getIdentifier();
           return ExprError();
Index: cfe/trunk/test/SemaCUDA/implicit-copy.cu
===================================================================
--- cfe/trunk/test/SemaCUDA/implicit-copy.cu
+++ cfe/trunk/test/SemaCUDA/implicit-copy.cu
@@ -0,0 +1,51 @@
+// RUN: %clang_cc1 -std=gnu++11 -triple nvptx64-unknown-unknown -fsyntax-only -verify %s
+// RUN: %clang_cc1 -std=gnu++11 -triple nvptx64-unknown-unknown -fcuda-is-device -fsyntax-only -verify %s
+
+struct CopyableH {
+  const CopyableH& operator=(const CopyableH& x) { return *this; }
+};
+struct CopyableD {
+  __attribute__((device)) const CopyableD& operator=(const CopyableD x) { return *this; }
+};
+
+struct SimpleH {
+  CopyableH b;
+};
+// expected-note@-3 2 {{candidate function (the implicit copy assignment operator) not viable: call to __host__ function from __device__ function}}
+// expected-note@-4 2 {{candidate function (the implicit move assignment operator) not viable: call to __host__ function from __device__ function}}
+
+struct SimpleD {
+  CopyableD b;
+};
+// expected-note@-3 2 {{candidate function (the implicit copy assignment operator) not viable: call to __device__ function from __host__ function}}
+// expected-note@-4 2 {{candidate function (the implicit move assignment operator) not viable: call to __device__ function from __host__ function}}
+
+void foo1hh() {
+  SimpleH a, b;
+  a = b;
+}
+__attribute__((device)) void foo1hd() {
+  SimpleH a, b;
+  a = b; // expected-error {{no viable overloaded}}
+}
+void foo1dh() {
+  SimpleD a, b;
+  a = b; // expected-error {{no viable overloaded}}
+}
+__attribute__((device)) void foo1dd() {
+  SimpleD a, b;
+  a = b;
+}
+
+void foo2hh(SimpleH &a, SimpleH &b) {
+  a = b;
+}
+__attribute__((device)) void foo2hd(SimpleH &a, SimpleH &b) {
+  a = b; // expected-error {{no viable overloaded}}
+}
+void foo2dh(SimpleD &a, SimpleD &b) {
+  a = b; // expected-error {{no viable overloaded}}
+}
+__attribute__((device)) void foo2dd(SimpleD &a, SimpleD &b) {
+  a = b;
+}
Index: cfe/trunk/test/SemaCUDA/implicit-member-target.cu
===================================================================
--- cfe/trunk/test/SemaCUDA/implicit-member-target.cu
+++ cfe/trunk/test/SemaCUDA/implicit-member-target.cu
@@ -146,11 +146,12 @@
 struct B7_with_copy_assign : A7_with_copy_assign {
 };
 
-// expected-note@-3 {{copy assignment operator of 'B7_with_copy_assign' is implicitly deleted}}
+// expected-note@-3 {{candidate function (the implicit copy assignment operator) not viable: call to __device__ function from __host__ function}}
+// expected-note@-4 {{candidate function (the implicit move assignment operator) not viable: call to __device__ function from __host__ function}}
 
 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}}
+  b1 = b2; // expected-error {{no viable overloaded '='}}
 }
 
 //------------------------------------------------------------------------------
@@ -176,9 +177,10 @@
 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}}
+// expected-note@-3 {{candidate function (the implicit copy assignment operator) not viable: call to __device__ function from __host__ function}}
+// expected-note@-4 {{candidate function (the implicit move assignment operator) not viable: call to __device__ function from __host__ function}}
 
 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}}
+  b1 = std::move(b2); // expected-error {{no viable overloaded '='}}
 }