Hi rnk, eliben,

In SemaCUDA all implicit functions were considered device host, this led to 
errors such as the following code snippet failing to compile:
```
struct Copyable {
  const Copyable& operator=(const Copyable& x) { return *this; }
};

struct Simple {
  Copyable b;
};

void foo() {
  Simple a, b;

  a = b;
}
```
Above the implicit copy assignment operator was inferred as host device but 
there was only a host assignment copy defined which is an error in device 
compilation mode.

http://reviews.llvm.org/D6565

Files:
  lib/Sema/SemaCUDA.cpp
  test/SemaCUDA/implicit-copy.cu
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;
+}
_______________________________________________
cfe-commits mailing list
[email protected]
http://lists.cs.uiuc.edu/mailman/listinfo/cfe-commits

Reply via email to