Consider calls from implict host device functions as valid in SemaCUDA

[jpienaar]

Summary

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.

[rnk]

lib/Sema/SemaCUDA.cpp

54 ↗

(On Diff #17041)

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?

test/SemaCUDA/implicit-copy.cu

16 ↗

(On Diff #17041)

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.

[jpienaar]

lib/Sema/SemaCUDA.cpp

54 ↗

(On Diff #17041)

Yes, although that might be too restrictive.

I was getting the following errors:

error: object of type 'Simple' cannot be assigned because its copy assignment operator is implicitly deleted
  a = b;
    ^
note: copy assignment operator of 'Simple' is implicitly deleted because field 'b' has no copy assignment operator
  Copyable b;
           ^

test/SemaCUDA/implicit-copy.cu

16 ↗

(On Diff #17041)

Yes, and it does that. The one where there is a problem is with

__device__ void foo(Simple &a, Simple &b) {
  a = b
}

There an error is flagged. In general, as CheckCUDATarget only looks at Caller and Callee, will we run into a problem in cases such as below

Host function calls a implicit function which calls a Device function

Now, a host function calling an implicit function is fine, an implicit function calling a device function is fine, but transitively we have a host function calling a device function. But this all seem to be just used for the method candidate set and errors are still reported if this happens. So it would seem that we could allow implicits in here while still catching such erroneous cases.

I have created a diff where implicits are considered always OK, I'll append it.

[jpienaar]

Changed the inference to treat all implicit functions uniformly. Extended test to include copy constructor.

[jpienaar]

Avoided adding an additional attribute.

[rnk]

lgtm

cfe/trunk/include/clang/Sema/Sema.h

enum CUDAFunctionTarget {
     CFT_Global,
     CFT_Host,
     CFT_HostDevice,
     CFT_InvalidTarget
   };
 
   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
   /// calls it needs to make to underlying base/field special members.
   /// \param ClassDecl the class for which the member is being created.
   /// \param CSM the kind of special member.
   /// \param MemberDecl the special member itself.
   /// \param ConstRHS true if this is a copy operation with a const object on

cfe/trunk/lib/Sema/SemaCUDA.cpp

if (D->hasAttr<CUDADeviceAttr>()) {
     return CFT_HostDevice;
   }
 
   return CFT_Host;
 }
 
 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)
     return true;
 
   // CUDA B.1.1 "The __device__ qualifier declares a function that is [...]
   // Callable from the device only."
   if (CallerTarget == CFT_Host && CalleeTarget == CFT_Device)
     return true;
 
   // CUDA B.1.2 "The __global__ qualifier declares a function that is [...]
   // Callable from the host only."
   // CUDA B.1.3 "The __host__ qualifier declares a function that is [...]
   // Callable from the host only."
   if ((CallerTarget == CFT_Device || CallerTarget == CFT_Global) &&
       (CalleeTarget == CFT_Host || CalleeTarget == CFT_Global))
     return true;
 
   // CUDA B.1.3 "The __device__ and __host__ qualifiers can be used together
   // 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;
   }
 
   return false;
 }
 

cfe/trunk/lib/Sema/SemaExpr.cpp

 ExprResult
 Sema::BuildDeclRefExpr(ValueDecl *D, QualType Ty, ExprValueKind VK,
                        const DeclarationNameInfo &NameInfo,
                        const CXXScopeSpec *SS, NamedDecl *FoundD,
                        const TemplateArgumentListInfo *TemplateArgs) {
   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();
         }
       }
 
   bool RefersToCapturedVariable =
       isa<VarDecl>(D) &&

cfe/trunk/test/SemaCUDA/implicit-copy.cu

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

cfe/trunk/test/SemaCUDA/implicit-member-target.cu

 struct A7_with_copy_assign {
   A7_with_copy_assign() {}
   __device__ A7_with_copy_assign& operator=(const 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 '='}}
 }
 
 //------------------------------------------------------------------------------
 // Test 8: move assignment operator
 
 // definitions for std::move
 namespace std {
 inline namespace foo {
struct A8_with_move_assign {
   A8_with_move_assign() {}
   __device__ A8_with_move_assign& operator=(A8_with_move_assign&&) {}
   __device__ A8_with_move_assign& operator=(const 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 '='}}
 }