Consider calls from implict host device functions as valid in SemaCUDA.

In SemaCUDA all implicit functions were considered host device, 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.

Differential Revision: http://reviews.llvm.org/D6565

llvm-svn: 224358

5 files changed, 66 insertions, 17 deletions

diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h
index 10d8b03fcc5..17154b636d0 100644
--- a/clang/include/clang/Sema/Sema.h
+++ b/clang/include/clang/Sema/Sema.h
@@ -8325,9 +8325,6 @@ public:
 
   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
diff --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp
index 46a83173fce..64222fbf8ac 100644
--- a/clang/lib/Sema/SemaCUDA.cpp
+++ b/clang/lib/Sema/SemaCUDA.cpp
@@ -62,12 +62,9 @@ Sema::CUDAFunctionTarget Sema::IdentifyCUDATarget(const FunctionDecl *D) {
 
 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 @@ bool Sema::CheckCUDATarget(CUDAFunctionTarget CallerTarget,
   // 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;
diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp
index 04497f3efd6..22cc925c25e 100644
--- a/clang/lib/Sema/SemaExpr.cpp
+++ b/clang/lib/Sema/SemaExpr.cpp
@@ -1590,11 +1590,10 @@ Sema::BuildDeclRefExpr(ValueDecl *D, QualType Ty, ExprValueKind VK,
   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();
diff --git a/clang/test/SemaCUDA/implicit-copy.cu b/clang/test/SemaCUDA/implicit-copy.cu
new file mode 100644
index 00000000000..b1b4887f561
--- /dev/null
+++ b/clang/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;
+}
diff --git a/clang/test/SemaCUDA/implicit-member-target.cu b/clang/test/SemaCUDA/implicit-member-target.cu
index 87b05e61a6f..6064560f2c6 100644
--- a/clang/test/SemaCUDA/implicit-member-target.cu
+++ b/clang/test/SemaCUDA/implicit-member-target.cu
@@ -146,11 +146,12 @@ struct A7_with_copy_assign {
 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 A8_with_move_assign {
 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 '='}}
 }