This revision was automatically updated to reflect the committed changes.
Closed by commit rL255911: [CUDA] Make vtable construction aware of host/device
side of CUDA compilation. (authored by tra).
Changed prior to commit:
http://reviews.llvm.org/D15309?vs=42341&id=43150#toc
Repository:
rL LLVM
http://reviews.llvm.org/D15309
Files:
cfe/trunk/lib/AST/RecordLayoutBuilder.cpp
cfe/trunk/lib/CodeGen/CGVTables.cpp
cfe/trunk/test/CodeGenCUDA/device-vtable.cu
Index: cfe/trunk/test/CodeGenCUDA/device-vtable.cu
===================================================================
--- cfe/trunk/test/CodeGenCUDA/device-vtable.cu
+++ cfe/trunk/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
+
Index: cfe/trunk/lib/CodeGen/CGVTables.cpp
===================================================================
--- cfe/trunk/lib/CodeGen/CGVTables.cpp
+++ cfe/trunk/lib/CodeGen/CGVTables.cpp
@@ -582,6 +582,24 @@
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) {
Index: cfe/trunk/lib/AST/RecordLayoutBuilder.cpp
===================================================================
--- cfe/trunk/lib/AST/RecordLayoutBuilder.cpp
+++ cfe/trunk/lib/AST/RecordLayoutBuilder.cpp
@@ -2025,6 +2025,21 @@
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.
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
