[clang] [Sema][CUDA][HIP] Reject RTTI operations in device code (PR #228469)
via cfe-commits
cfe-commits at lists.llvm.org
Fri Oct 2 07:53:57 PDT 2026
llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-clang
Author: Steffen Larsen (steffenlarsen)
<details>
<summary>Changes</summary>
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
---
Full diff: https://github.com/llvm/llvm-project/pull/228469.diff
4 Files Affected:
- (modified) clang/include/clang/Basic/DiagnosticSemaKinds.td (+4)
- (modified) clang/lib/Sema/SemaCast.cpp (+7)
- (modified) clang/lib/Sema/SemaExprCXX.cpp (+14)
- (added) clang/test/SemaCUDA/device-rtti.cu (+100)
``````````diff
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 at -1 {{cannot use 'dynamic_cast' in __device__ function as RTTI is not available in device code}}
+ (void)dynamic_cast<D &>(*b);
+ // expected-error at -1 {{cannot use 'dynamic_cast' in __device__ function as RTTI is not available in device code}}
+ (void)typeid(D);
+ // expected-error at -1 {{cannot use 'typeid' in __device__ function as RTTI is not available in device code}}
+ (void)typeid(*b);
+ // expected-error at -1 {{cannot use 'typeid' in __device__ function as RTTI is not available in device code}}
+ (void)sizeof(typeid(int));
+ // expected-error at -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 at -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 at -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 at -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 at -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 at -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 at -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 at -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 at -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 at -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(); }
``````````
</details>
https://github.com/llvm/llvm-project/pull/228469
More information about the cfe-commits
mailing list