[CUDA] emit vtables only for classes with methods usable on this side of compilation.

[tra]

Summary

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 be generated on
device side, because we are not going to generate any code for the
host-only methods during device-side compilation.

For classes that mix host and device virtual methods, nvcc generates vtable
with null pointers for methods that are not available on the current side.

• During key method calculation only considers methods valid for the current CUDA compilation. This prevents vtable generation for classes that have no virtual methods usable on this side of CUDA compilation.
• For vtables we emit during current CUDA compilation, entries for the methods that are not valid in this mode are set to null. This allows us to generate valid vtables for classes that have mix of host/device virtual methods.

[cfe-commits]

• msg-25712-156.txt159 BDownload

[tra]

Removed unnecessary temporary variable and 'else'.

[jpienaar]

Looks good to me.

lib/AST/RecordLayoutBuilder.cpp

2003 ↗

(On Diff #42122)

Add a comment to explain asymmetry. Perhaps a general comment at the start of this section explaining would be the most useful.

test/CodeGenCUDA/device-vtable.cu

36 ↗

(On Diff #42122)

What is the current behavior in this case? Should an error be reported?

[tra]

test/CodeGenCUDA/device-vtable.cu

36 ↗

(On Diff #42122)

Hmm. nvcc actually accepts mixed virtual methods and generates vtable with NULL pointers for methods that are not available on particular side. That shifts responsibility to end-user which is
reasonable in this situation. A warning about partial availability of some virtual method would be helpful here, I think.

OK. Let me try matching nvcc's behavior and add a warning.

[tra]

Changed key method algorithm to ignore methods that we will not emit during this compilation (previous version didn't allow mixed host/device virtual methods).

For vtables with mixed host/device virtual methods, use NULL for methods we're not going to codegen.

[tra]

Ping.
@rsmith -- Richard, can you take a look at this change as well?

[jpienaar]

Looks good, thanks

test/CodeGenCUDA/device-vtable.cu

37 ↗

(On Diff #42341)

Does the comment need to be updated to reflect vtable with NULL pointers?

cfe/trunk/lib/AST/RecordLayoutBuilder.cpp

for (const CXXMethodDecl *MD : RD->methods()) {
 
     // In certain ABIs, ignore functions with out-of-line inline definitions.
     if (!allowInlineFunctions) {
       const FunctionDecl *Def;
       if (MD->hasBody(Def) && Def->isInlineSpecified())
         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.
     if (MD->hasAttr<DLLImportAttr>() && !RD->hasAttr<DLLImportAttr>())
       return nullptr;
 
     // We found it.
     return MD;

cfe/trunk/lib/CodeGen/CGVTables.cpp

case VTableComponent::CK_DeletingDtorPointer: {
       case VTableComponent::CK_CompleteDtorPointer:
         GD = GlobalDecl(Component.getDestructorDecl(), Dtor_Complete);
         break;
       case VTableComponent::CK_DeletingDtorPointer:
         GD = GlobalDecl(Component.getDestructorDecl(), Dtor_Deleting);
         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) {
           llvm::FunctionType *Ty = 
             llvm::FunctionType::get(CGM.VoidTy, /*isVarArg=*/false);
           StringRef PureCallName = CGM.getCXXABI().GetPureVirtualCallName();
           PureVirtualFn = CGM.CreateRuntimeFunction(Ty, PureCallName);
           PureVirtualFn = llvm::ConstantExpr::getBitCast(PureVirtualFn,

cfe/trunk/test/CodeGenCUDA/device-vtable.cu

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