Author: yaxunl Date: Tue Jun 5 08:11:02 2018 New Revision: 334021 URL: http://llvm.org/viewvc/llvm-project?rev=334021&view=rev Log: [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 Modified: cfe/trunk/lib/CodeGen/CodeGenModule.cpp cfe/trunk/test/CodeGenCUDA/device-vtable.cu Modified: cfe/trunk/lib/CodeGen/CodeGenModule.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CodeGenModule.cpp?rev=334021&r1=334020&r2=334021&view=diff ============================================================================== --- cfe/trunk/lib/CodeGen/CodeGenModule.cpp (original) +++ cfe/trunk/lib/CodeGen/CodeGenModule.cpp Tue Jun 5 08:11:02 2018 @@ -4900,7 +4900,7 @@ llvm::Constant *CodeGenModule::GetAddrOf // 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() && Modified: cfe/trunk/test/CodeGenCUDA/device-vtable.cu URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenCUDA/device-vtable.cu?rev=334021&r1=334020&r2=334021&view=diff ============================================================================== --- cfe/trunk/test/CodeGenCUDA/device-vtable.cu (original) +++ cfe/trunk/test/CodeGenCUDA/device-vtable.cu Tue Jun 5 08:11:02 2018 @@ -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 _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
