Index: include/clang/Basic/Attr.td
===================================================================
--- include/clang/Basic/Attr.td
+++ include/clang/Basic/Attr.td
@@ -563,6 +563,13 @@
   let Documentation = [Undocumented];
 }
 
+def CUDAImplicitTarget : InheritableAttr {
+  let Spellings = [];
+  let Subjects = SubjectList<[Function]>;
+  let LangOpts = [CUDA];
+  let Documentation = [Undocumented];
+}
+
 def CUDAInvalidTarget : InheritableAttr {
   let Spellings = [];
   let Subjects = SubjectList<[Function]>;
Index: include/clang/Basic/DiagnosticSemaKinds.td
===================================================================
--- include/clang/Basic/DiagnosticSemaKinds.td
+++ include/clang/Basic/DiagnosticSemaKinds.td
@@ -3068,8 +3068,8 @@
     "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__|invalid}1 function from"
-    " %select{__device__|__global__|__host__|__host__ __device__|invalid}2 function">;
+    "%select{__device__|__global__|__host__|__host__ __device__|implicit|invalid}1 function from"
+    " %select{__device__|__global__|__host__|__host__ __device__|implicit|invalid}2 function">;
 def note_implicit_member_target_infer_collision : Note<
     "implicit %select{"
     "default constructor|"
Index: include/clang/Sema/Sema.h
===================================================================
--- include/clang/Sema/Sema.h
+++ include/clang/Sema/Sema.h
@@ -8295,6 +8295,7 @@
     CFT_Global,
     CFT_Host,
     CFT_HostDevice,
+    CFT_ImplicitTarget,
     CFT_InvalidTarget
   };
 
Index: lib/Sema/SemaCUDA.cpp
===================================================================
--- lib/Sema/SemaCUDA.cpp
+++ lib/Sema/SemaCUDA.cpp
@@ -52,9 +52,7 @@
   } else if (D->hasAttr<CUDAHostAttr>()) {
     return CFT_Host;
   } else if (D->isImplicit()) {
-    // Some implicit declarations (like intrinsic functions) are not marked.
-    // Set the most lenient target on them for maximal flexibility.
-    return CFT_HostDevice;
+    return CFT_ImplicitTarget;
   }
 
   return CFT_Host;
@@ -73,6 +71,10 @@
   if (CallerTarget == CFT_InvalidTarget || CalleeTarget == CFT_InvalidTarget)
     return true;
 
+  // If one of the targets is implicit, the check always passes.
+  if (CallerTarget == CFT_ImplicitTarget || CalleeTarget == CFT_ImplicitTarget)
+    return false;
+
   // CUDA B.1.1 "The __device__ qualifier declares a function that is [...]
   // Callable from the device only."
   if (CallerTarget == CFT_Host && CalleeTarget == CFT_Device)
Index: test/SemaCUDA/implicit-copy.cu
===================================================================
--- test/SemaCUDA/implicit-copy.cu
+++ test/SemaCUDA/implicit-copy.cu
@@ -0,0 +1,46 @@
+// 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 {  // expected-note 4 {{candidate function}}
+  CopyableH b;
+};
+struct SimpleD {  // expected-note 4 {{candidate function}}
+  CopyableD b;
+};
+
+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: test/SemaCUDA/implicit-member-target.cu
===================================================================
--- test/SemaCUDA/implicit-member-target.cu
+++ 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 '='}}
 }