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

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

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;

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();

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;
+}

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 '='}}
 }