dnrops
/
clang
mirror of https://github.com/microsoft/clang.git
1
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity

[CUDA][HIP] Do not emit type info when compiling for device 

CUDA/HIP does not support RTTI on device side, therefore there
is no point of emitting type info when compiling for device.

Emitting type info for device not only clutters the IR with useless
global variables, but also causes undefined symbol at linking
since vtable for cxxabiv1::class_type_info has external linkage.

Differential Revision: https://reviews.llvm.org/D47694


git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@334021 91177308-0d34-0410-b5e6-96231b3b80d8
This commit is contained in:
Yaxun Liu 2018-06-05 15:11:02 +00:00
parent 2cf50762e0
commit 571a2502f2
2 changed files with 10 additions and 3 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
lib/CodeGen/CodeGenModule.cpp
View File

@ -4900,7 +4900,7 @@ llvm::Constant *CodeGenModule::GetAddrOfRTTIDescriptor(QualType Ty,
// Return a bogus pointer if RTTI is disabled, unless it's for EH.
// FIXME: should we even be calling this method if RTTI is disabled
// and it's not for EH?
if (!ForEH && !getLangOpts().RTTI)
if ((!ForEH && !getLangOpts().RTTI) || getLangOpts().CUDAIsDevice)
return llvm::Constant::getNullValue(Int8PtrTy);
if (ForEH && Ty->isObjCObjectPointerType() &&

11
test/CodeGenCUDA/device-vtable.cu
View File

@ -19,7 +19,9 @@ struct H {
//CHECK-HOST: @_ZTV1H =
//CHECK-HOST-SAME: @_ZN1H6methodEv
//CHECK-DEVICE-NOT: @_ZTV1H =
//CHECK-DEVICE-NOT: @_ZTVN10__cxxabiv117__class_type_infoE
//CHECK-DEVICE-NOT: @_ZTS1H
//CHECK-DEVICE-NOT: @_ZTI1H
struct D {
__device__ virtual void method();
};
@ -27,7 +29,9 @@ struct D {
//CHECK-DEVICE: @_ZTV1D
//CHECK-DEVICE-SAME: @_ZN1D6methodEv
//CHECK-HOST-NOT: @_ZTV1D
//CHECK-DEVICE-NOT: @_ZTVN10__cxxabiv117__class_type_infoE
//CHECK-DEVICE-NOT: @_ZTS1D
//CHECK-DEVICE-NOT: @_ZTI1D
// 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
@ -45,6 +49,9 @@ struct HD {
// CHECK-HOST-NOT: @_ZN2HD8d_methodEv
// CHECK-HOST-SAME: null
// CHECK-BOTH-SAME: ]
// CHECK-DEVICE-NOT: @_ZTVN10__cxxabiv117__class_type_infoE
// CHECK-DEVICE-NOT: @_ZTS2HD
// CHECK-DEVICE-NOT: @_ZTI2HD
void H::method() {}
//CHECK-HOST: define void @_ZN1H6methodEv