[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

diff --git a/lib/CodeGen/CodeGenModule.cpp b/lib/CodeGen/CodeGenModule.cpp
index 3098e0f984c2fb6c59a9fc527a3de9b65cab866c..b57004f7e122a25e2253dcdbb04435f1f32a3524 100644 (file)
--- a/lib/CodeGen/CodeGenModule.cpp
+++ b/lib/CodeGen/CodeGenModule.cpp
@@ -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() &&

diff --git a/test/CodeGenCUDA/device-vtable.cu b/test/CodeGenCUDA/device-vtable.cu
index 9730e404caa4fc7c1caa5de8e6f54aa0cd808534..a7307dcdaf9bdbea80c6449bad72057968713bd6 100644 (file)
--- a/test/CodeGenCUDA/device-vtable.cu
+++ b/test/CodeGenCUDA/device-vtable.cu
@@ -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