[CUDA] Make vtable construction aware of host/device side of CUDA compilation.

C++ emits vtables for classes that have key function present in the
current TU. While we compile CUDA the fact that key function was found
in this TU does not mean that we are going to generate code for it. E.g.
vtable for a class with host-only methods should not (and can not) be
generated on device side, because we'll never generate code for them
during device-side compilation.

This patch adds an extra CUDA-specific check during key method computation
and filters out potential key methods that are not suitable for this side
of CUDA compilation.

When we codegen vtable, entries for unsuitable methods are set to null.

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

git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@255911 91177308-0d34-0410-b5e6-96231b3b80d8

diff --git a/lib/AST/RecordLayoutBuilder.cpp b/lib/AST/RecordLayoutBuilder.cpp
index 0c76838a6d21b6d89bee2604283acbd9b2796b66..bc3c2a831c47f1259d5c0f21c876fe360474d060 100644 (file)
--- a/lib/AST/RecordLayoutBuilder.cpp
+++ b/lib/AST/RecordLayoutBuilder.cpp
@@ -2025,6 +2025,21 @@ static const CXXMethodDecl *computeKeyFunction(ASTContext &Context,
         continue;
     }
 
+    if (Context.getLangOpts().CUDA) {
+      // While compiler may see key method in this TU, during CUDA
+      // compilation we should ignore methods that are not accessible
+      // on this side of compilation.
+      if (Context.getLangOpts().CUDAIsDevice) {
+        // In device mode ignore methods without __device__ attribute.
+        if (!MD->hasAttr<CUDADeviceAttr>())
+          continue;
+      } else {
+        // In host mode ignore __device__-only methods.
+        if (!MD->hasAttr<CUDAHostAttr>() && MD->hasAttr<CUDADeviceAttr>())
+          continue;
+      }
+    }
+
     // If the key function is dllimport but the class isn't, then the class has
     // no key function. The DLL that exports the key function won't export the
     // vtable in this case.

diff --git a/lib/CodeGen/CGVTables.cpp b/lib/CodeGen/CGVTables.cpp
index 797c4085c19ac93e003e7b7900423d5004c629f3..c8f3add67762965b67339a4763b20a63c9e9584b 100644 (file)
--- a/lib/CodeGen/CGVTables.cpp
+++ b/lib/CodeGen/CGVTables.cpp
@@ -582,6 +582,24 @@ llvm::Constant *CodeGenVTables::CreateVTableInitializer(
         break;
       }
 
+      if (CGM.getLangOpts().CUDA) {
+        // Emit NULL for methods we can't codegen on this
+        // side. Otherwise we'd end up with vtable with unresolved
+        // references.
+        const CXXMethodDecl *MD = cast<CXXMethodDecl>(GD.getDecl());
+        // OK on device side: functions w/ __device__ attribute
+        // OK on host side: anything except __device__-only functions.
+        bool CanEmitMethod = CGM.getLangOpts().CUDAIsDevice
+                                 ? MD->hasAttr<CUDADeviceAttr>()
+                                 : (MD->hasAttr<CUDAHostAttr>() ||
+                                    !MD->hasAttr<CUDADeviceAttr>());
+        if (!CanEmitMethod) {
+          Init = llvm::ConstantExpr::getNullValue(Int8PtrTy);
+          break;
+        }
+        // Method is acceptable, continue processing as usual.
+      }
+
       if (cast<CXXMethodDecl>(GD.getDecl())->isPure()) {
         // We have a pure virtual member function.
         if (!PureVirtualFn) {

diff --git a/test/CodeGenCUDA/device-vtable.cu b/test/CodeGenCUDA/device-vtable.cu
new file mode 100644 (file)
index 0000000..9730e40
--- /dev/null
+++ b/test/CodeGenCUDA/device-vtable.cu
@@ -0,0 +1,61 @@
+// REQUIRES: x86-registered-target
+// REQUIRES: nvptx-registered-target
+
+// Make sure we don't emit vtables for classes with methods that have
+// inappropriate target attributes. Currently it's mostly needed in
+// order to avoid emitting vtables for host-only classes on device
+// side where we can't codegen them.
+
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm -o - %s \
+// RUN:     | FileCheck %s -check-prefix=CHECK-HOST -check-prefix=CHECK-BOTH
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -emit-llvm -o - %s \
+// RUN:     | FileCheck %s -check-prefix=CHECK-DEVICE -check-prefix=CHECK-BOTH
+
+#include "Inputs/cuda.h"
+
+struct H  {
+  virtual void method();
+};
+//CHECK-HOST: @_ZTV1H =
+//CHECK-HOST-SAME: @_ZN1H6methodEv
+//CHECK-DEVICE-NOT: @_ZTV1H =
+
+struct D  {
+   __device__ virtual void method();
+};
+
+//CHECK-DEVICE: @_ZTV1D
+//CHECK-DEVICE-SAME: @_ZN1D6methodEv
+//CHECK-HOST-NOT: @_ZTV1D
+
+// This is the case with mixed host and device virtual methods.  It's
+// impossible to emit a valid vtable in that case because only host or
+// only device methods would be available during host or device
+// compilation. At the moment Clang (and NVCC) emit NULL pointers for
+// unavailable methods,
+struct HD  {
+  virtual void h_method();
+  __device__ virtual void d_method();
+};
+// CHECK-BOTH: @_ZTV2HD
+// CHECK-DEVICE-NOT: @_ZN2HD8h_methodEv
+// CHECK-DEVICE-SAME: null
+// CHECK-DEVICE-SAME: @_ZN2HD8d_methodEv
+// CHECK-HOST-SAME: @_ZN2HD8h_methodEv
+// CHECK-HOST-NOT: @_ZN2HD8d_methodEv
+// CHECK-HOST-SAME: null
+// CHECK-BOTH-SAME: ]
+
+void H::method() {}
+//CHECK-HOST: define void @_ZN1H6methodEv
+
+void __device__ D::method() {}
+//CHECK-DEVICE: define void @_ZN1D6methodEv
+
+void __device__ HD::d_method() {}
+// CHECK-DEVICE: define void @_ZN2HD8d_methodEv
+// CHECK-HOST-NOT: define void @_ZN2HD8d_methodEv
+void HD::h_method() {}
+// CHECK-HOST: define void @_ZN2HD8h_methodEv
+// CHECK-DEVICE-NOT: define void @_ZN2HD8h_methodEv
+