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

[David Blaikie via cfe-commits]

On Mon, Dec 7, 2015 at 2:44 PM, Artem Belevich via cfe-commits <
cfe-commits at lists.llvm.org> wrote:

> 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;
>

Guess you intended to use this local variable, but didn't? ^


> +    for (const CXXMethodDecl *MD : RD->methods())
> +      if (Context.getLangOpts().CUDAIsDevice &&
> !MD->hasAttr<CUDADeviceAttr>())
> +        return nullptr;
> +      else if (!Context.getLangOpts().CUDAIsDevice &&
>

Drop the else after return (just "if/return/if/return")
(and/or you could roll both conditions into one if test - potentially even
using "std::any_of" (or llvm::any_of))


> +               !MD->hasAttr<CUDAHostAttr>() &&
> MD->hasAttr<CUDADeviceAttr>())
> +        return nullptr;
> +  }
> +
>    for (const CXXMethodDecl *MD : RD->methods()) {
>      if (!MD->isVirtual())
>        continue;
>
>
>
> _______________________________________________
> cfe-commits mailing list
> cfe-commits at lists.llvm.org
> http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
>
>
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20151207/c8285185/attachment-0001.html>