On Mon, Dec 7, 2015 at 2:44 PM, Artem Belevich via cfe-commits < cfe-commits@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@lists.llvm.org > http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits > >
_______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
