https://github.com/steffenlarsen created https://github.com/llvm/llvm-project/pull/228469
Currently RTTI operations, i.e. dynamic_cast and typeid, are silently ignored in device code. This does not match the behavior of NVCC, which rejects these operations in device code. Make Sema reject these operations in device code, emitting a diagnostic. An effect of this is that the compiler may reject code that previously compiled, such as code that used RTTI in dead device code that the compiler would remove before reaching the linker. However, cases would currently fail if optimizations are disabled. Assisted-by: Claude Opus 5.5 >From d3827a3c22f27e363dfdb9e28bc53fa72ab028dc Mon Sep 17 00:00:00 2001 From: Steffen Holst Larsen <[email protected]> Date: Fri, 2 Oct 2026 09:19:36 -0500 Subject: [PATCH] [Sema][CUDA][HIP] Reject RTTI operations in device code Currently RTTI operations, i.e. dynamic_cast and typeid, are silently ignored in device code. This does not match the behavior of NVCC, which rejects these operations in device code. Make Sema reject these operations in device code, emitting a diagnostic. An effect of this is that the compiler may reject code that previously compiled, such as code that used RTTI in dead device code that the compiler would remove before reaching the linker. However, cases would currently fail if optimizations are disabled. Assisted-by: Claude Opus 5.5 Signed-off-by: Steffen Holst Larsen <[email protected]> --- .../clang/Basic/DiagnosticSemaKinds.td | 4 + clang/lib/Sema/SemaCast.cpp | 7 ++ clang/lib/Sema/SemaExprCXX.cpp | 14 +++ clang/test/SemaCUDA/device-rtti.cu | 100 ++++++++++++++++++ 4 files changed, 125 insertions(+) create mode 100644 clang/test/SemaCUDA/device-rtti.cu diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 9208aba1445d7eb..0b59df6bc529f8d 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -9817,6 +9817,10 @@ def note_cuda_conflicting_device_function_declared_here : Note< def err_cuda_device_exceptions : Error< "cannot use '%0' in " "%select{__device__|__global__|__host__|__host__ __device__}1 function">; +def err_cuda_device_rtti : Error< + "cannot use '%0' in " + "%select{__device__|__global__|__host__|__host__ __device__}1 function as " + "RTTI is not available in device code">; def err_dynamic_var_init : Error< "dynamic initialization is not supported for " "__device__, __constant__, __shared__, and __managed__ variables">; diff --git a/clang/lib/Sema/SemaCast.cpp b/clang/lib/Sema/SemaCast.cpp index c797bdb11e0cbdc..6f3debf498db5b0 100644 --- a/clang/lib/Sema/SemaCast.cpp +++ b/clang/lib/Sema/SemaCast.cpp @@ -24,6 +24,7 @@ #include "clang/Lex/Preprocessor.h" #include "clang/Sema/Initialization.h" #include "clang/Sema/SemaAMDGPU.h" +#include "clang/Sema/SemaCUDA.h" #include "clang/Sema/SemaHLSL.h" #include "clang/Sema/SemaObjC.h" #include "clang/Sema/SemaRISCV.h" @@ -987,6 +988,12 @@ void CastOperation::CheckDynamicCast() { return; } + // Similarly, dynamic_cast is not available in CUDA device code, except for + // dynamic_cast to void*. + if (Self.getLangOpts().CUDA && !DestPointee->isVoidType()) + Self.CUDA().DiagIfDeviceCode(OpRange.getBegin(), diag::err_cuda_device_rtti) + << "dynamic_cast" << Self.CUDA().CurrentTarget(); + // Warns when dynamic_cast is used with RTTI data disabled. if (!Self.getLangOpts().RTTIData) { bool MicrosoftABI = diff --git a/clang/lib/Sema/SemaExprCXX.cpp b/clang/lib/Sema/SemaExprCXX.cpp index b89c97f2b8900a7..85f312639f8ba47 100644 --- a/clang/lib/Sema/SemaExprCXX.cpp +++ b/clang/lib/Sema/SemaExprCXX.cpp @@ -537,10 +537,22 @@ bool Sema::checkLiteralOperatorId(const CXXScopeSpec &SS, llvm_unreachable("unknown nested name specifier kind"); } +/// RTTI is not available in CUDA/HIP device code, so typeid can't be used +/// there. A dependent operand is checked once the template is instantiated. +static void diagnoseCUDADeviceTypeid(Sema &S, SourceLocation TypeidLoc, + bool IsDependent) { + if (S.getLangOpts().CUDA && !IsDependent) + S.CUDA().DiagIfDeviceCode(TypeidLoc, diag::err_cuda_device_rtti) + << "typeid" << S.CUDA().CurrentTarget(); +} + ExprResult Sema::BuildCXXTypeId(QualType TypeInfoType, SourceLocation TypeidLoc, TypeSourceInfo *Operand, SourceLocation RParenLoc) { + diagnoseCUDADeviceTypeid(*this, TypeidLoc, + Operand->getType()->isDependentType()); + // C++ [expr.typeid]p4: // The top-level cv-qualifiers of the lvalue expression or the type-id // that is the operand of typeid are always ignored. @@ -568,6 +580,8 @@ ExprResult Sema::BuildCXXTypeId(QualType TypeInfoType, SourceLocation TypeidLoc, Expr *E, SourceLocation RParenLoc) { + diagnoseCUDADeviceTypeid(*this, TypeidLoc, E && E->isTypeDependent()); + bool WasEvaluated = false; if (E && !E->isTypeDependent()) { if (E->hasPlaceholderType()) { diff --git a/clang/test/SemaCUDA/device-rtti.cu b/clang/test/SemaCUDA/device-rtti.cu new file mode 100644 index 000000000000000..ce7c13c13b6eb90 --- /dev/null +++ b/clang/test/SemaCUDA/device-rtti.cu @@ -0,0 +1,100 @@ +// RUN: %clang_cc1 -fcuda-is-device -fsyntax-only -verify=expected,dev %s +// RUN: %clang_cc1 -fsyntax-only -verify %s +// RUN: %clang_cc1 -x hip -fcuda-is-device -fsyntax-only -verify=expected,dev %s +// RUN: %clang_cc1 -x hip -fsyntax-only -verify %s + +#include "Inputs/cuda.h" + +namespace std { +class type_info {}; +} // namespace std + +struct B { + __host__ __device__ virtual ~B() {} +}; +struct D : B {}; + +void host(B *b) { + (void)dynamic_cast<D *>(b); + (void)typeid(*b); +} + +__device__ void device(B *b, D *d) { + (void)dynamic_cast<D *>(b); + // expected-error@-1 {{cannot use 'dynamic_cast' in __device__ function as RTTI is not available in device code}} + (void)dynamic_cast<D &>(*b); + // expected-error@-1 {{cannot use 'dynamic_cast' in __device__ function as RTTI is not available in device code}} + (void)typeid(D); + // expected-error@-1 {{cannot use 'typeid' in __device__ function as RTTI is not available in device code}} + (void)typeid(*b); + // expected-error@-1 {{cannot use 'typeid' in __device__ function as RTTI is not available in device code}} + (void)sizeof(typeid(int)); + // expected-error@-1 {{cannot use 'typeid' in __device__ function as RTTI is not available in device code}} + + // As with -fno-rtti, these don't use RTTI and are allowed. + (void)dynamic_cast<void *>(b); + (void)dynamic_cast<B *>(d); +} + +__global__ void kernel(B *b) { + (void)typeid(*b); + // expected-error@-1 {{cannot use 'typeid' in __global__ function as RTTI is not available in device code}} +} + +// Check that it's an error to use RTTI from a __host__ __device__ function if +// and only if it's codegen'ed for device. + +__host__ __device__ void hd1(B *b) { + (void)dynamic_cast<D *>(b); + // dev-error@-1 {{cannot use 'dynamic_cast' in __host__ __device__ function as RTTI is not available in device code}} +} + +// No error, never instantiated on device. +inline __host__ __device__ void hd2(B *b) { (void)typeid(*b); } +void call_hd2(B *b) { hd2(b); } + +// Error, instantiated on device. +inline __host__ __device__ void hd3(B *b) { + (void)typeid(*b); + // dev-error@-1 {{cannot use 'typeid' in __host__ __device__ function as RTTI is not available in device code}} +} +__device__ void call_hd3(B *b) { hd3(b); } +// dev-note@-1 {{called by 'call_hd3'}} + +// Templates are checked when they are instantiated. +template <class T> __device__ T *tmpl_unused(B *b) { + return dynamic_cast<T *>(b); +} + +template <class T> __device__ T *tmpl(B *b) { + return dynamic_cast<T *>(b); + // expected-error@-1 {{cannot use 'dynamic_cast' in __device__ function as RTTI is not available in device code}} +} +__device__ void call_tmpl(B *b) { tmpl<D>(b); } +// expected-note@-1 {{in instantiation of function template specialization 'tmpl<D>' requested here}} + +template <class T> __device__ void tmpl_typeid(T *t) { + (void)typeid(*t); + // expected-error@-1 {{cannot use 'typeid' in __device__ function as RTTI is not available in device code}} +} +__device__ void call_tmpl_typeid(B *b) { tmpl_typeid(b); } +// expected-note@-1 {{in instantiation of function template specialization 'tmpl_typeid<B>' requested here}} + +template <class T> __device__ void tmpl_typeid_unused(T *t) { + (void)typeid(*t); +} + +// A non-dependent operand is checked once, in the template definition. +template <class T> __device__ void tmpl_nondependent() { + (void)typeid(int); + // expected-error@-1 {{cannot use 'typeid' in __device__ function as RTTI is not available in device code}} +} +__device__ void call_tmpl_nondependent() { tmpl_nondependent<int>(); } + +// A host virtual function in a class that also has device virtual functions +// is not device code. +struct Fix { + __device__ virtual void run() {} + virtual D *init(B *b) { return dynamic_cast<D *>(b); } +}; +__device__ void use_fix() { Fix f; f.run(); } _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
