yaxunl updated this revision to Diff 149892.
yaxunl added a comment.
Revised by John's comments.
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,7 @@
// 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() &&
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,7 @@
// 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() &&
_______________________________________________
cfe-commits mailing list
[email protected]
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
