Index: lib/Sema/SemaCUDA.cpp
===================================================================
--- lib/Sema/SemaCUDA.cpp
+++ lib/Sema/SemaCUDA.cpp
@@ -51,9 +51,9 @@
     return CFT_Device;
   } 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.
+  } else if (D->isImplicit() && D->getBuiltinID()) {
+    // Intrinsic functions cannot (at present) be marked as __device__ so set
+    // the most lenient target on them for maximal flexibility.
     return CFT_HostDevice;
   }
 
Index: test/SemaCUDA/implicit-copy.cu
===================================================================
--- test/SemaCUDA/implicit-copy.cu
+++ test/SemaCUDA/implicit-copy.cu
@@ -0,0 +1,16 @@
+// RUN: %clang_cc1 -std=gnu++11 -triple nvptx64-unknown-unknown -fcuda-is-device -fsyntax-only -verify %s
+
+struct Copyable {
+  const Copyable& operator=(const Copyable& x) { return *this; }
+};
+
+struct Simple {
+  Copyable b;
+};
+
+// expected-no-diagnostics
+void foo() {
+  Simple a, b;
+
+  a = b;
+}