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/4] [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 9208aba1445d7..0b59df6bc529f 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 c797bdb11e0cb..6f3debf498db5 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 b89c97f2b8900..85f312639f8ba 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 0000000000000..ce7c13c13b6eb --- /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/4] 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 ce7c13c13b6eb..6af154c17d71b 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/4] 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 0b59df6bc529f..485b2be352d21 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 bf990f21e7b63..51474a0017b35 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 090030ea82503..714793c6bdbed 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 6f3debf498db5..cea9a593e58ad 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 785c7409d8952..c29b7e48fc8a3 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 85f312639f8ba..577923aa31cea 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 6af154c17d71b..f1503a65db85d 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/4] 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 f1503a65db85d..d85a28e4cb245 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}} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
