tra created this revision.
tra added reviewers: rsmith, jingyue, jpienaar.
tra added a subscriber: cfe-commits.
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.
During CUDA compilation this patch checks virtual methods' target attributes
and
returns key function only if all virtual methods in the class are suitable for
the current
compilation mode.
http://reviews.llvm.org/D15309
Files:
lib/AST/RecordLayoutBuilder.cpp
test/CodeGenCUDA/device-vtable.cu
Index: test/CodeGenCUDA/device-vtable.cu
===================================================================
--- /dev/null
+++ test/CodeGenCUDA/device-vtable.cu
@@ -0,0 +1,55 @@
+// 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
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -emit-llvm -o
- %s \
+// RUN: | FileCheck %s -check-prefix=CHECK-DEVICE
+
+#include "Inputs/cuda.h"
+
+class H {
+ public:
+ virtual void method();
+};
+//CHECK-HOST: @_ZTV1H =
+//CHECK-HOST-SAME: @_ZN1H6methodEv
+//CHECK-DEVICE-NOT: @_ZTV1H =
+
+class D {
+ public:
+ __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. For now we'll not emit such vtable at all.
+class HD {
+ public:
+ virtual void h_method();
+ __device__ virtual void d_method();
+};
+//CHECK-BOTH-NOT: @_ZTV2HD
+
+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: lib/AST/RecordLayoutBuilder.cpp
===================================================================
--- lib/AST/RecordLayoutBuilder.cpp
+++ lib/AST/RecordLayoutBuilder.cpp
@@ -1996,6 +1996,16 @@
bool allowInlineFunctions =
Context.getTargetInfo().getCXXABI().canKeyFunctionBeInline();
+ if (Context.getLangOpts().CUDA) {
+ const bool IsDevice = Context.getLangOpts().CUDAIsDevice;
+ for (const CXXMethodDecl *MD : RD->methods())
+ if (Context.getLangOpts().CUDAIsDevice && !MD->hasAttr<CUDADeviceAttr>())
+ return nullptr;
+ else if (!Context.getLangOpts().CUDAIsDevice &&
+ !MD->hasAttr<CUDAHostAttr>() && MD->hasAttr<CUDADeviceAttr>())
+ return nullptr;
+ }
+
for (const CXXMethodDecl *MD : RD->methods()) {
if (!MD->isVirtual())
continue;
Index: test/CodeGenCUDA/device-vtable.cu
===================================================================
--- /dev/null
+++ test/CodeGenCUDA/device-vtable.cu
@@ -0,0 +1,55 @@
+// 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
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -emit-llvm -o - %s \
+// RUN: | FileCheck %s -check-prefix=CHECK-DEVICE
+
+#include "Inputs/cuda.h"
+
+class H {
+ public:
+ virtual void method();
+};
+//CHECK-HOST: @_ZTV1H =
+//CHECK-HOST-SAME: @_ZN1H6methodEv
+//CHECK-DEVICE-NOT: @_ZTV1H =
+
+class D {
+ public:
+ __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. For now we'll not emit such vtable at all.
+class HD {
+ public:
+ virtual void h_method();
+ __device__ virtual void d_method();
+};
+//CHECK-BOTH-NOT: @_ZTV2HD
+
+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: lib/AST/RecordLayoutBuilder.cpp
===================================================================
--- lib/AST/RecordLayoutBuilder.cpp
+++ lib/AST/RecordLayoutBuilder.cpp
@@ -1996,6 +1996,16 @@
bool allowInlineFunctions =
Context.getTargetInfo().getCXXABI().canKeyFunctionBeInline();
+ if (Context.getLangOpts().CUDA) {
+ const bool IsDevice = Context.getLangOpts().CUDAIsDevice;
+ for (const CXXMethodDecl *MD : RD->methods())
+ if (Context.getLangOpts().CUDAIsDevice && !MD->hasAttr<CUDADeviceAttr>())
+ return nullptr;
+ else if (!Context.getLangOpts().CUDAIsDevice &&
+ !MD->hasAttr<CUDAHostAttr>() && MD->hasAttr<CUDADeviceAttr>())
+ return nullptr;
+ }
+
for (const CXXMethodDecl *MD : RD->methods()) {
if (!MD->isVirtual())
continue;
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
