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

[yaxunl]

Summary

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.

[rjmccall]

Why not just have the driver disable RTTI in the frontend invocation?

[yaxunl]

In D47694#1120367, @rjmccall wrote:

Why not just have the driver disable RTTI in the frontend invocation?

CUDA/HIP uses single source for device and host. The host code may depend on RTTI,
e.g., an application may include some boost headers which will fail if RTTI is disabled,
therefore RTTI cannot be disabled when compiling device code.

[rjmccall]

In D47694#1120375, @yaxunl wrote:

In D47694#1120367, @rjmccall wrote:

Why not just have the driver disable RTTI in the frontend invocation?

CUDA/HIP uses single source for device and host. The host code may depend on RTTI,
e.g., an application may include some boost headers which will fail if RTTI is disabled,
therefore RTTI cannot be disabled when compiling device code.

It's a single source file, yes, but behind the scenes you do a separate invocation of the compiler to re-parse that file for device code-generation, right?

[yaxunl]

In D47694#1121037, @rjmccall wrote:

In D47694#1120375, @yaxunl wrote:

In D47694#1120367, @rjmccall wrote:

Why not just have the driver disable RTTI in the frontend invocation?

CUDA/HIP uses single source for device and host. The host code may depend on RTTI,
e.g., an application may include some boost headers which will fail if RTTI is disabled,
therefore RTTI cannot be disabled when compiling device code.

It's a single source file, yes, but behind the scenes you do a separate invocation of the compiler to re-parse that file for device code-generation, right?

That's true. However, when compiling for device, the source file is still the original source code which are for both host and device. When compiling
for device, the compiler still needs to parse the source code for host, and there is no way to disable RTTI only for the device code.

[rjmccall]

Oh, I see, because you're worried that the host code might contain dynamic_cast or similar features which would complain if RTTI were disabled.

getLangOpts().CUDAIsDevice implies getLangOpts().CUDA, so I think you can just check the former. Otherwise LGTM.

[yaxunl]

Revised by John's comments.

[rjmccall]

Thanks, LGTM.

lib/CodeGen/CodeGenModule.cpp

   return llvm::ConstantStruct::getAnon(Fields);
 }
 
 llvm::Constant *CodeGenModule::GetAddrOfRTTIDescriptor(QualType Ty,
                                                        bool ForEH) {
   // 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().CUDA && getLangOpts().CUDAIsDevice))
     return llvm::Constant::getNullValue(Int8PtrTy);
 
   if (ForEH && Ty->isObjCObjectPointerType() &&
       LangOpts.ObjCRuntime.isGNUFamily())
     return ObjCRuntime->GetEHType(Ty);
 
   return getCXXABI().getAddrOfRTTIDescriptor(Ty);
 }

test/CodeGenCUDA/device-vtable.cu

 #include "Inputs/cuda.h"
 
 struct H  {
   virtual void method();
 };
 //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();
 };
 
 //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
 // compilation. At the moment Clang (and NVCC) emit NULL pointers for
 // unavailable methods,
 struct HD  {
   virtual void h_method();
   __device__ virtual void d_method();
 };
 // CHECK-BOTH: @_ZTV2HD
 // CHECK-DEVICE-NOT: @_ZN2HD8h_methodEv
 // CHECK-DEVICE-SAME: null
 // CHECK-DEVICE-SAME: @_ZN2HD8d_methodEv
 // CHECK-HOST-SAME: @_ZN2HD8h_methodEv
 // 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
 
 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