https://github.com/steffenlarsen updated https://github.com/llvm/llvm-project/pull/228469
>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 1/6] [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 9208aba1445d7e..0b59df6bc529f8 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 c797bdb11e0cbd..6f3debf498db5b 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 b89c97f2b8900a..85f312639f8ba4 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 00000000000000..ce7c13c13b6eb9 --- /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(); } >From dfee9a610ed050e4052d33ff4cfc965175235cbd Mon Sep 17 00:00:00 2001 From: Steffen Holst Larsen <[email protected]> Date: Mon, 5 Oct 2026 02:04:13 -0500 Subject: [PATCH 2/6] Extend testing Signed-off-by: Steffen Holst Larsen <[email protected]> --- clang/test/SemaCUDA/device-rtti.cu | 43 ++++++++++++++++++++++++++++++ 1 file changed, 43 insertions(+) diff --git a/clang/test/SemaCUDA/device-rtti.cu b/clang/test/SemaCUDA/device-rtti.cu index ce7c13c13b6eb9..6af154c17d71b0 100644 --- a/clang/test/SemaCUDA/device-rtti.cu +++ b/clang/test/SemaCUDA/device-rtti.cu @@ -19,6 +19,11 @@ void host(B *b) { (void)typeid(*b); } +__host__ void explicit_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}} @@ -30,6 +35,11 @@ __device__ void device(B *b, D *d) { // 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}} + (void)sizeof(dynamic_cast<D *>(b)); + // expected-error@-1 {{cannot use 'dynamic_cast' in __device__ function as RTTI is not available in device code}} + decltype(dynamic_cast<D *>(b)) ptr = d; + // expected-error@-1 {{cannot use 'dynamic_cast' in __device__ function as RTTI is not available in device code}} + (void)ptr; // As with -fno-rtti, these don't use RTTI and are allowed. (void)dynamic_cast<void *>(b); @@ -37,10 +47,29 @@ __device__ void device(B *b, D *d) { } __global__ void kernel(B *b) { + (void)dynamic_cast<D *>(b); + // expected-error@-1 {{cannot use 'dynamic_cast' in __global__ function as RTTI is not available in device code}} (void)typeid(*b); // expected-error@-1 {{cannot use 'typeid' in __global__ function as RTTI is not available in device code}} } +struct S { + __device__ D *member(B *b) { return dynamic_cast<D *>(b); } + // expected-error@-1 {{cannot use 'dynamic_cast' in __device__ function as RTTI is not available in device code}} +}; + +// A __device__ lambda is device code, and an unannotated lambda is +// __host__ __device__, so it is device code when called from device code. +__device__ void lambdas(B *b) { + auto dev = [] __device__ (B *p) { return dynamic_cast<D *>(p); }; + // expected-error@-1 {{cannot use 'dynamic_cast' in __device__ function as RTTI is not available in device code}} + auto hd = [](B *p) { return dynamic_cast<D *>(p); }; + // dev-error@-1 {{cannot use 'dynamic_cast' in __host__ __device__ function as RTTI is not available in device code}} + dev(b); + hd(b); + // dev-note@-1 {{called by 'lambdas'}} +} + // Check that it's an error to use RTTI from a __host__ __device__ function if // and only if it's codegen'ed for device. @@ -91,6 +120,20 @@ template <class T> __device__ void tmpl_nondependent() { } __device__ void call_tmpl_nondependent() { tmpl_nondependent<int>(); } +template <class T> __device__ T *explicit_inst(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}} +} +template __device__ D *explicit_inst<D>(B *); +// expected-note@-1 {{in instantiation of function template specialization 'explicit_inst<D>' requested here}} + +template <class T> __global__ void kernel_tmpl(B *b) { + (void)typeid(*static_cast<T *>(b)); + // expected-error@-1 {{cannot use 'typeid' in __global__ function as RTTI is not available in device code}} +} +template __global__ void kernel_tmpl<D>(B *); +// expected-note@-1 {{in instantiation of function template specialization 'kernel_tmpl<D>' requested here}} + // A host virtual function in a class that also has device virtual functions // is not device code. struct Fix { >From 4c8548a5ddb1d032ea09d1af1f01344f47638b1c Mon Sep 17 00:00:00 2001 From: Steffen Holst Larsen <[email protected]> Date: Mon, 5 Oct 2026 02:54:09 -0500 Subject: [PATCH 3/6] RTTI outside function bodies Signed-off-by: Steffen Holst Larsen <[email protected]> --- .../clang/Basic/DiagnosticSemaKinds.td | 5 + clang/include/clang/Sema/SemaCUDA.h | 17 +++ clang/lib/Sema/SemaCUDA.cpp | 89 +++++++++++++ clang/lib/Sema/SemaCast.cpp | 3 +- clang/lib/Sema/SemaExpr.cpp | 7 ++ clang/lib/Sema/SemaExprCXX.cpp | 3 +- clang/test/SemaCUDA/device-rtti.cu | 118 ++++++++++++++++++ 7 files changed, 238 insertions(+), 4 deletions(-) diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 0b59df6bc529f8..485b2be352d213 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -9821,6 +9821,11 @@ 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_cuda_device_rtti_var_init : Error< + "cannot use '%0' in the initializer of a device variable as RTTI is not " + "available in device code">; +def note_cuda_device_rtti_default_used_here : Note< + "%select{default member initializer|default argument}0 used here">; def err_dynamic_var_init : Error< "dynamic initialization is not supported for " "__device__, __constant__, __shared__, and __managed__ variables">; diff --git a/clang/include/clang/Sema/SemaCUDA.h b/clang/include/clang/Sema/SemaCUDA.h index bf990f21e7b63a..51474a0017b35a 100644 --- a/clang/include/clang/Sema/SemaCUDA.h +++ b/clang/include/clang/Sema/SemaCUDA.h @@ -107,6 +107,15 @@ class SemaCUDA : public SemaBase { /// Same as DiagIfDeviceCode, with "host" and "device" switched. SemaDiagnosticBuilder DiagIfHostCode(SourceLocation Loc, unsigned DiagID); + /// Diagnoses a use of RTTI at \p Loc if it is in device function or the + /// initializer of a device variable, where RTTI is not available. + void checkRTTIUse(SourceLocation Loc, StringRef Op); + + /// Diagnoses a use of RTTI in \p Init, the default argument or default + /// member initializer of \p D if it is used at \p UseLoc in device code. + void checkRTTIInDefaultInit(const ValueDecl *D, const Expr *Init, + SourceLocation UseLoc); + /// Determines whether the given function is a CUDA device/host/kernel/etc. /// function. /// @@ -290,6 +299,14 @@ class SemaCUDA : public SemaBase { private: unsigned ForceHostDeviceDepth = 0; + /// Whether the current context is the initializer of a device variable. + bool isDeviceVarInit() const; + + /// The default arguments and default member initializers, paired with the + /// function using them, that checkRTTIInDefaultInit has checked. + llvm::DenseSet<std::pair<const ValueDecl *, const FunctionDecl *>> + RTTICheckedDefaultInits; + friend class ASTReader; friend class ASTWriter; }; diff --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp index 090030ea82503a..714793c6bdbed9 100644 --- a/clang/lib/Sema/SemaCUDA.cpp +++ b/clang/lib/Sema/SemaCUDA.cpp @@ -13,6 +13,7 @@ #include "clang/Sema/SemaCUDA.h" #include "clang/AST/ASTContext.h" #include "clang/AST/Decl.h" +#include "clang/AST/DynamicRecursiveASTVisitor.h" #include "clang/AST/EvaluatedExprVisitor.h" #include "clang/AST/ExprCXX.h" #include "clang/Basic/Cuda.h" @@ -1015,6 +1016,94 @@ Sema::SemaDiagnosticBuilder SemaCUDA::DiagIfHostCode(SourceLocation Loc, return SemaDiagnosticBuilder(DiagKind, Loc, DiagID, CurFunContext, SemaRef); } +bool SemaCUDA::isDeviceVarInit() const { + return CurCUDATargetCtx.Kind == CTCK_InitGlobalVar && + CurCUDATargetCtx.Target == CUDAFunctionTarget::Device; +} + +void SemaCUDA::checkRTTIUse(SourceLocation Loc, StringRef Op) { + assert(getLangOpts().CUDA && "Should only be called during CUDA compilation"); + // A default argument is checked where it is used. + if (isa_and_present<ParmVarDecl>( + SemaRef.currentEvaluationContext().ManglingContextDecl)) + return; + + if (SemaRef.getCurFunctionDecl(/*AllowLambda=*/true)) { + DiagIfDeviceCode(Loc, diag::err_cuda_device_rtti) << Op << CurrentTarget(); + return; + } + + // Outside a function, only the initializer of a device variable is device + // code. A default member initializer is checked where it is used. + if (isDeviceVarInit()) + Diag(Loc, diag::err_cuda_device_rtti_var_init) << Op; +} + +namespace { +/// Finds a use of RTTI, as diagnosed by checkRTTIUse, in a default argument or +/// default member initializer. +struct RTTIUseFinder : DynamicRecursiveASTVisitor { + const Expr *Use = nullptr; + + RTTIUseFinder() { + // Look into nested default arguments and default member initializers. + ShouldVisitImplicitCode = true; + // Lambda bodies are checked as functions of their own. + ShouldVisitLambdaBody = false; + } + + bool VisitCXXTypeidExpr(CXXTypeidExpr *E) override { + Use = E; + return false; + } + + bool VisitCXXDynamicCastExpr(CXXDynamicCastExpr *E) override { + // As with -fno-rtti, upcasts and dynamic_cast to void* don't use RTTI. + QualType DestTy = E->getType(); + if (E->getCastKind() != CK_Dynamic || + (DestTy->isPointerType() && DestTy->getPointeeType()->isVoidType())) + return true; + Use = E; + return false; + } +}; +} // namespace + +void SemaCUDA::checkRTTIInDefaultInit(const ValueDecl *D, const Expr *Init, + SourceLocation UseLoc) { + assert(getLangOpts().CUDA && "Should only be called during CUDA compilation"); + assert((isa<FieldDecl, ParmVarDecl>(D)) && + "Expected a default member initializer or default argument"); + const FunctionDecl *CurFn = SemaRef.getCurFunctionDecl(/*AllowLambda=*/true); + if (!CurFn && !isDeviceVarInit()) + return; + RTTIUseFinder Finder; + Finder.TraverseStmt(const_cast<Expr *>(Init)); + const Expr *Use = Finder.Use; + if (!Use) + return; + // An initializer can be built more than once for the same use. + if (!RTTICheckedDefaultInits.insert({D, CurFn}).second) + return; + + StringRef Op = isa<CXXTypeidExpr>(Use) ? "typeid" : "dynamic_cast"; + bool IsDefaultArg = isa<ParmVarDecl>(D); + if (!CurFn) { + Diag(Use->getBeginLoc(), diag::err_cuda_device_rtti_var_init) << Op; + Diag(UseLoc, diag::note_cuda_device_rtti_default_used_here) << IsDefaultArg; + return; + } + { + SemaDiagnosticBuilder DB = + DiagIfDeviceCode(Use->getBeginLoc(), diag::err_cuda_device_rtti); + DB << Op << CurrentTarget(); + // Let the note below follow the error, as SemaBase::Diag does. + SemaRef.IsLastErrorImmediate = DB.isImmediate(); + } + DiagIfDeviceCode(UseLoc, diag::note_cuda_device_rtti_default_used_here) + << IsDefaultArg; +} + bool SemaCUDA::CheckCall(SourceLocation Loc, FunctionDecl *Callee) { assert(getLangOpts().CUDA && "Should only be called during CUDA compilation"); assert(Callee && "Callee may not be null."); diff --git a/clang/lib/Sema/SemaCast.cpp b/clang/lib/Sema/SemaCast.cpp index 6f3debf498db5b..cea9a593e58ada 100644 --- a/clang/lib/Sema/SemaCast.cpp +++ b/clang/lib/Sema/SemaCast.cpp @@ -991,8 +991,7 @@ void CastOperation::CheckDynamicCast() { // 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(); + Self.CUDA().checkRTTIUse(OpRange.getBegin(), "dynamic_cast"); // Warns when dynamic_cast is used with RTTI data disabled. if (!Self.getLangOpts().RTTIData) { diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp index 785c7409d89525..c29b7e48fc8a3b 100644 --- a/clang/lib/Sema/SemaExpr.cpp +++ b/clang/lib/Sema/SemaExpr.cpp @@ -5900,6 +5900,10 @@ ExprResult Sema::BuildCXXDefaultArgExpr(SourceLocation CallLoc, /*SkipImmediateInvocations=*/NestedDefaultChecking)) return ExprError(); + if (getLangOpts().CUDA) + CUDA().checkRTTIInDefaultInit(Param, Init ? Init : Param->getDefaultArg(), + CallLoc); + return CXXDefaultArgExpr::Create(Context, InitializationContext->Loc, Param, Init, InitializationContext->Context); } @@ -5969,6 +5973,9 @@ ExprResult Sema::BuildCXXDefaultInitInternal(SourceLocation Loc, return ExprError(); } + if (getLangOpts().CUDA) + CUDA().checkRTTIInDefaultInit(Field, InClassInit, Loc); + // CWG2631 // An immediate invocation that is not evaluated where it appears is // evaluated and checked for whether it is a constant expression at the diff --git a/clang/lib/Sema/SemaExprCXX.cpp b/clang/lib/Sema/SemaExprCXX.cpp index 85f312639f8ba4..577923aa31ceab 100644 --- a/clang/lib/Sema/SemaExprCXX.cpp +++ b/clang/lib/Sema/SemaExprCXX.cpp @@ -542,8 +542,7 @@ bool Sema::checkLiteralOperatorId(const CXXScopeSpec &SS, 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(); + S.CUDA().checkRTTIUse(TypeidLoc, "typeid"); } ExprResult Sema::BuildCXXTypeId(QualType TypeInfoType, diff --git a/clang/test/SemaCUDA/device-rtti.cu b/clang/test/SemaCUDA/device-rtti.cu index 6af154c17d71b0..f1503a65db85df 100644 --- a/clang/test/SemaCUDA/device-rtti.cu +++ b/clang/test/SemaCUDA/device-rtti.cu @@ -141,3 +141,121 @@ struct Fix { virtual D *init(B *b) { return dynamic_cast<D *>(b); } }; __device__ void use_fix() { Fix f; f.run(); } + +// Outside a function, the initializer of a device variable is device code. +__device__ const std::type_info *device_var = &typeid(int); +// expected-error@-1 {{cannot use 'typeid' in the initializer of a device variable as RTTI is not available in device code}} +__constant__ const std::type_info *constant_var = &typeid(int); +// expected-error@-1 {{cannot use 'typeid' in the initializer of a device variable as RTTI is not available in device code}} +const std::type_info *host_var = &typeid(int); + +// A default member initializer is checked where it is used. +struct MemberInit { + const std::type_info *t = &typeid(int); +}; +// expected-error@#member_init {{cannot use 'typeid' in __device__ function as RTTI is not available in device code}} +// dev-error@#member_init {{cannot use 'typeid' in __host__ __device__ function as RTTI is not available in device code}} +// dev-note@#member_init_struct {{default member initializer used here}} + +__device__ void member_init_ctor() { + MemberInit m; + // dev-note@-1 {{called by 'member_init_ctor'}} + (void)m; +} +__device__ void member_init_aggregate() { + MemberInit m{}; + // expected-note@-1 {{default member initializer used here}} + (void)m; +} +void member_init_host() { + MemberInit m; + MemberInit n{}; + (void)m; + (void)n; +} + +struct MemberInitDeviceCtor { + B *b = nullptr; + D *d = dynamic_cast<D *>(b); + // expected-error@-1 {{cannot use 'dynamic_cast' in __device__ function as RTTI is not available in device code}} + __device__ MemberInitDeviceCtor() {} + // expected-note@-1 {{default member initializer used here}} +}; + +// No error, the constructor is never used on device. +struct MemberInitHDCtor { + const std::type_info *t = &typeid(int); + __host__ __device__ MemberInitHDCtor() {} +}; +void member_init_hd_ctor_host() { MemberInitHDCtor m; } + +template <class T> struct MemberInitTmpl { // #member_init_tmpl + const std::type_info *t = &typeid(T); + // dev-error@-1 {{cannot use 'typeid' in __host__ __device__ function as RTTI is not available in device code}} +}; +// dev-note@#member_init_tmpl {{default member initializer used here}} +__device__ void member_init_tmpl() { + MemberInitTmpl<int> m; + // dev-note@-1 {{called by 'member_init_tmpl'}} + (void)m; +} + +// A default argument is checked where it is used. +__device__ void default_arg(const std::type_info *t = &typeid(int)); // #default_arg +// expected-error@-1 {{cannot use 'typeid' in __device__ function as RTTI is not available in device code}} +__device__ void use_default_arg() { + default_arg(); + // expected-note@-1 {{default argument used here}} +} + +inline __host__ __device__ void hd_default_arg(const std::type_info *t = &typeid(int)) {} +// expected-error@-1 {{cannot use 'typeid' in __device__ function as RTTI is not available in device code}} +void use_hd_default_arg_host() { hd_default_arg(); } +__device__ void use_hd_default_arg_device() { + hd_default_arg(); + // expected-note@-1 {{default argument used here}} +} + +// Diagnosed once, at the use, although the default argument is parsed in a +// __device__ function. +__device__ void local_default_arg() { + __device__ void local(const std::type_info *t = &typeid(int)); + // expected-error@-1 {{cannot use 'typeid' in __device__ function as RTTI is not available in device code}} + local(); + // expected-note@-1 {{default argument used here}} +} + +// A default argument used in a default member initializer. +struct NestedDefaultArg { // #nested_struct + const std::type_info *t = (default_arg(), nullptr); +}; +// dev-error@#default_arg {{cannot use 'typeid' in __host__ __device__ function as RTTI is not available in device code}} +// dev-note@#nested_struct {{default member initializer used here}} +__device__ void nested_default_arg() { + NestedDefaultArg n; + // dev-note@-1 {{called by 'nested_default_arg'}} + (void)n; +} + +// A lambda body is checked as a function of its own, while the initializer of +// a capture is part of the default member initializer. +struct LambdaBody { + const std::type_info *t = [] { return &typeid(int); }(); // #lambda_body +}; +// dev-error@#lambda_body {{cannot use 'typeid' in __host__ __device__ function as RTTI is not available in device code}} +// dev-note@#lambda_body {{called by 'lambda_body'}} +__device__ void lambda_body() { + LambdaBody l; + (void)l; +} + +struct LambdaCapture { // #lambda_capture_struct + const std::type_info *t = [p = &typeid(int)] { return p; }(); + // dev-error@-1 {{cannot use 'typeid' in __host__ __device__ function as RTTI is not available in device code}} +}; +// dev-note@#lambda_capture_struct {{default member initializer used here}} +__device__ void lambda_capture() { + LambdaCapture l; + // dev-note@-1 {{called by 'lambda_capture'}} + (void)l; +} >From d202f4dc6788d815c335b0682fd154a6691602d4 Mon Sep 17 00:00:00 2001 From: Steffen Holst Larsen <[email protected]> Date: Mon, 5 Oct 2026 03:50:30 -0500 Subject: [PATCH 4/6] Add missing markers Signed-off-by: Steffen Holst Larsen <[email protected]> --- clang/test/SemaCUDA/device-rtti.cu | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/clang/test/SemaCUDA/device-rtti.cu b/clang/test/SemaCUDA/device-rtti.cu index f1503a65db85df..d85a28e4cb2459 100644 --- a/clang/test/SemaCUDA/device-rtti.cu +++ b/clang/test/SemaCUDA/device-rtti.cu @@ -150,8 +150,8 @@ __constant__ const std::type_info *constant_var = &typeid(int); const std::type_info *host_var = &typeid(int); // A default member initializer is checked where it is used. -struct MemberInit { - const std::type_info *t = &typeid(int); +struct MemberInit { // #member_init_struct + const std::type_info *t = &typeid(int); // #member_init }; // expected-error@#member_init {{cannot use 'typeid' in __device__ function as RTTI is not available in device code}} // dev-error@#member_init {{cannot use 'typeid' in __host__ __device__ function as RTTI is not available in device code}} >From 2dc927323e8a1c540e149558a1eac49abcd71346 Mon Sep 17 00:00:00 2001 From: Steffen Holst Larsen <[email protected]> Date: Tue, 6 Oct 2026 02:51:47 -0500 Subject: [PATCH 5/6] Switch to ConstDynamicRecursiveASTVisitor Signed-off-by: Steffen Holst Larsen <[email protected]> --- clang/lib/Sema/SemaCUDA.cpp | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp index 714793c6bdbed9..3be7e977522c3c 100644 --- a/clang/lib/Sema/SemaCUDA.cpp +++ b/clang/lib/Sema/SemaCUDA.cpp @@ -1042,7 +1042,7 @@ void SemaCUDA::checkRTTIUse(SourceLocation Loc, StringRef Op) { namespace { /// Finds a use of RTTI, as diagnosed by checkRTTIUse, in a default argument or /// default member initializer. -struct RTTIUseFinder : DynamicRecursiveASTVisitor { +struct RTTIUseFinder : ConstDynamicRecursiveASTVisitor { const Expr *Use = nullptr; RTTIUseFinder() { @@ -1052,12 +1052,12 @@ struct RTTIUseFinder : DynamicRecursiveASTVisitor { ShouldVisitLambdaBody = false; } - bool VisitCXXTypeidExpr(CXXTypeidExpr *E) override { + bool VisitCXXTypeidExpr(const CXXTypeidExpr *E) override { Use = E; return false; } - bool VisitCXXDynamicCastExpr(CXXDynamicCastExpr *E) override { + bool VisitCXXDynamicCastExpr(const CXXDynamicCastExpr *E) override { // As with -fno-rtti, upcasts and dynamic_cast to void* don't use RTTI. QualType DestTy = E->getType(); if (E->getCastKind() != CK_Dynamic || @@ -1078,7 +1078,7 @@ void SemaCUDA::checkRTTIInDefaultInit(const ValueDecl *D, const Expr *Init, if (!CurFn && !isDeviceVarInit()) return; RTTIUseFinder Finder; - Finder.TraverseStmt(const_cast<Expr *>(Init)); + Finder.TraverseStmt(Init); const Expr *Use = Finder.Use; if (!Use) return; >From 81a3ffb217e7285390e8e75bae278b6be22439d4 Mon Sep 17 00:00:00 2001 From: Steffen Holst Larsen <[email protected]> Date: Tue, 6 Oct 2026 03:40:52 -0500 Subject: [PATCH 6/6] Add dependent default arguments and default member initialiser test cases Signed-off-by: Steffen Holst Larsen <[email protected]> --- clang/test/SemaCUDA/device-rtti.cu | 91 ++++++++++++++++++++++++++++++ 1 file changed, 91 insertions(+) diff --git a/clang/test/SemaCUDA/device-rtti.cu b/clang/test/SemaCUDA/device-rtti.cu index d85a28e4cb2459..941cfd54fe90a9 100644 --- a/clang/test/SemaCUDA/device-rtti.cu +++ b/clang/test/SemaCUDA/device-rtti.cu @@ -200,6 +200,31 @@ __device__ void member_init_tmpl() { (void)m; } +// Whether a dependent default member initializer uses RTTI can depend on the +// template arguments. +template <class T> struct MemberInitDepCast { // #member_init_dep_cast + B *b = nullptr; + T *t = dynamic_cast<T *>(b); + // dev-error@-1 {{cannot use 'dynamic_cast' in __host__ __device__ function as RTTI is not available in device code}} +}; +// dev-note@#member_init_dep_cast {{default member initializer used here}} +__device__ void member_init_dep_cast() { + MemberInitDepCast<D> d; + // dev-note@-1 {{called by 'member_init_dep_cast'}} + MemberInitDepCast<B> b; + (void)d; + (void)b; +} + +// No error, only instantiated for host code. +template <class T> struct MemberInitTmplHost { + const std::type_info *t = &typeid(T); +}; +void member_init_tmpl_host() { + MemberInitTmplHost<int> m; + (void)m; +} + // A default argument is checked where it is used. __device__ void default_arg(const std::type_info *t = &typeid(int)); // #default_arg // expected-error@-1 {{cannot use 'typeid' in __device__ function as RTTI is not available in device code}} @@ -225,6 +250,60 @@ __device__ void local_default_arg() { // expected-note@-1 {{default argument used here}} } +// A dependent default argument is checked when it is instantiated for a use. +template <class T> +__device__ void dep_default_arg(const std::type_info *t = &typeid(T)); // #dep_default_arg +// expected-error@#dep_default_arg {{cannot use 'typeid' in __device__ function as RTTI is not available in device code}} +__device__ void use_dep_default_arg() { + dep_default_arg<int>(); + // expected-note@-1 {{default argument used here}} + // No error, the default argument is not used. + dep_default_arg<float>(nullptr); +} + +// Whether a dependent default argument uses RTTI can depend on the template +// arguments. +template <class T> +__device__ void dep_cast_default_arg(T *p = dynamic_cast<T *>(static_cast<B *>(nullptr))); +// expected-error@-1 {{cannot use 'dynamic_cast' in __device__ function as RTTI is not available in device code}} +__device__ void use_dep_cast_default_arg() { + dep_cast_default_arg<D>(); + // expected-note@-1 {{default argument used here}} + dep_cast_default_arg<B>(); + dep_cast_default_arg<void>(); +} + +template <class T> struct DepDefaultArgMember { + __device__ void f(const std::type_info *t = &typeid(T)); + // expected-error@-1 {{cannot use 'typeid' in __device__ function as RTTI is not available in device code}} +}; +__device__ void use_dep_default_arg_member() { + DepDefaultArgMember<int>().f(); + // expected-note@-1 {{default argument used here}} +} + +template <class T> +inline __host__ __device__ void hd_dep_default_arg(const std::type_info *t = &typeid(T)) {} +// expected-error@-1 {{cannot use 'typeid' in __device__ function as RTTI is not available in device code}} +void use_hd_dep_default_arg_host() { hd_dep_default_arg<int>(); } +__device__ void use_hd_dep_default_arg_device() { + hd_dep_default_arg<float>(); + // expected-note@-1 {{default argument used here}} +} + +// A non-dependent default argument of a template is checked for each +// specialization that uses it. +template <class T> +__device__ void nondep_default_arg(T, const std::type_info *t = &typeid(int)); // #nondep_default_arg +// expected-error@#nondep_default_arg 2{{cannot use 'typeid' in __device__ function as RTTI is not available in device code}} +__device__ void use_nondep_default_arg() { + nondep_default_arg(1); + // expected-note@-1 {{default argument used here}} + nondep_default_arg(1.0); + // expected-note@-1 {{default argument used here}} + nondep_default_arg(2); +} + // A default argument used in a default member initializer. struct NestedDefaultArg { // #nested_struct const std::type_info *t = (default_arg(), nullptr); @@ -237,6 +316,18 @@ __device__ void nested_default_arg() { (void)n; } +// A dependent default argument used in a dependent default member initializer. +template <class T> struct DepNestedDefaultArg { // #dep_nested_struct + const std::type_info *t = (dep_default_arg<T>(), nullptr); +}; +// dev-error@#dep_default_arg {{cannot use 'typeid' in __host__ __device__ function as RTTI is not available in device code}} +// dev-note@#dep_nested_struct {{default member initializer used here}} +__device__ void dep_nested_default_arg() { + DepNestedDefaultArg<long> n; + // dev-note@-1 {{called by 'dep_nested_default_arg'}} + (void)n; +} + // A lambda body is checked as a function of its own, while the initializer of // a capture is part of the default member initializer. struct LambdaBody { _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
