yaxunl created this revision.
yaxunl added reviewers: rjmccall, tra.
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.
https://reviews.llvm.org/D47694
Files:
lib/CodeGen/CodeGenModule.cpp
test/CodeGenCUDA/device-vtable.cu
Index: test/CodeGenCUDA/device-vtable.cu
===================================================================
--- test/CodeGenCUDA/device-vtable.cu
+++ test/CodeGenCUDA/device-vtable.cu
@@ -19,15 +19,19 @@
//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
@@ -45,6 +49,9 @@
// 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
Index: lib/CodeGen/CodeGenModule.cpp
===================================================================
--- lib/CodeGen/CodeGenModule.cpp
+++ lib/CodeGen/CodeGenModule.cpp
@@ -4899,7 +4899,8 @@
// 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() &&
Index: test/CodeGenCUDA/device-vtable.cu
===================================================================
--- test/CodeGenCUDA/device-vtable.cu
+++ test/CodeGenCUDA/device-vtable.cu
@@ -19,15 +19,19 @@
//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
@@ -45,6 +49,9 @@
// 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
Index: lib/CodeGen/CodeGenModule.cpp
===================================================================
--- lib/CodeGen/CodeGenModule.cpp
+++ lib/CodeGen/CodeGenModule.cpp
@@ -4899,7 +4899,8 @@
// 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() &&
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
