[LLVMbugs] [Bug 20495] New: Clang incorrectly marks implicitly-defined constructors as __host__ __device__ in CUDA mode

[bugzilla-daemon at llvm.org]

http://llvm.org/bugs/show_bug.cgi?id=20495

            Bug ID: 20495
           Summary: Clang incorrectly marks implicitly-defined
                    constructors as __host__ __device__ in CUDA mode
           Product: clang
           Version: trunk
          Hardware: PC
                OS: Linux
            Status: NEW
          Severity: normal
          Priority: P
         Component: Frontend
          Assignee: unassignedclangbugs at nondot.org
          Reporter: eliben at gmail.com
                CC: llvmbugs at cs.uiuc.edu
    Classification: Unclassified

Consider this code:

  struct A {
   A() {aa = 12;}

   int aa;
  };

  struct B : A {
    int bb;
  };

  void foo() {
   B b;
  }

Clang currently gives an error when compiling in -x cuda mode:

pc.cc:12:4: error: call to implicitly-deleted default constructor of 'B'
 B b;
   ^
pc.cc:7:12: note: default constructor of 'B' is implicitly deleted because base
class 'A' has no default constructor
struct B : A {
           ^
1 error generated.

----

Another example which is a bit more telling:

  struct A { A(const A&); };
  struct B : A {};

  extern B b1;
  B b2 = std::move(b1);

Clang complains:

h.cu:3:8: error: no matching constructor for initialization of 'A'
struct B : A {};
       ^
h.cu:2:12: note: candidate constructor not viable: call to __host__ function
from __host__ __device__ function
struct A { A(const A&); };
           ^
h.cu:6:8: note: implicit move constructor for 'B' first required here
B b2 = std::move(b1);
       ^


In both cases the errors should not happen, but they do because Clang marks
implicitly created constructors with "__host__ __device__", so when they call
__host__ constructors of subclasses, it detects an error.

This recognition happens in Sema::IdentifyCUDATarget:

Sema::CUDAFunctionTarget Sema::IdentifyCUDATarget(const FunctionDecl *D) {
  // Implicitly declared functions (e.g. copy constructors) are
  // __host__ __device__
  if (D->isImplicit())
    return CFT_HostDevice;
 ...
 ...

However, this is incorrect. Marking a function __host__ __device__ means that
it must not invoke code that is __host__ only or __device__ only. So this is
too restrictive. Since the user did not actually specify the constructor and it
was inferred implicitly, this constructor should be available both in host and
device code. This doesn't mean it's __host__ __device__. It means it can call
both host and device code - these are different things.

-- 
You are receiving this mail because:
You are on the CC list for the bug.
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/llvm-bugs/attachments/20140730/6ae5a60b/attachment.html>