[PATCH] Consider only implicit builtin functions as device host in SemaCUDA

[Reid Kleckner]

================
Comment at: lib/Sema/SemaCUDA.cpp:54
@@ -53,4 +53,3 @@
     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
----------------
Builtins are for things like __builtin_popcount(), but your test case is looking at implicitly defined C++ special members, right? What was happening on the test case you have previously?

================
Comment at: test/SemaCUDA/implicit-copy.cu:16
@@ +15,2 @@
+  a = b;
+}
----------------
This seems like another interesting test case:

  struct Copyable {
    __device__ const Copyable& operator=(const Copyable& x) { return *this; }
  };
  struct Simple {
    Copyable b;
  };
  void foo(Simple &a, Simple &b) {
    a = b;
  }

Simple's implicit copy ctor should be rejected when the host is the target because it calls a device-only method from the host, right?

I'm also curious what happens when someone uses FP math builtins like __builtin_cos that might have a reasonable lowering on a GPU.

http://reviews.llvm.org/D6565