[clang] [CIR] Handle dynamic_cast and typeid without RTTI in device code (PR #228031)
Steffen Larsen via cfe-commits
cfe-commits at lists.llvm.org
Thu Oct 1 03:50:59 PDT 2026
https://github.com/steffenlarsen created https://github.com/llvm/llvm-project/pull/228031
RTTI is not emitted for CUDA/HIP device code. As in OGCG, the RTTI attribute returns a null pointer instead of a reference to the RTTI descriptor, but dynamic_cast and typeid expected a GlobalViewAttr and crashed on it.
Make the RTTI parameters of cir.dyn_cast_info optional. CIRGen leaves them out when RTTI is not available, and the Itanium ABI lowering passes null pointers to __dynamic_cast instead. The verifier requires either both or neither, and since the attribute alias is named after the RTTI, such an attribute is printed without one. typeid returns the null pointer, like OGCG.
>From 623c6699b043f13aca15e6cbf14cd3c046c0939a Mon Sep 17 00:00:00 2001
From: Steffen Holst Larsen <sholstla at amd.com>
Date: Tue, 22 Sep 2026 08:42:53 -0500
Subject: [PATCH] [CIR] Handle dynamic_cast and typeid without RTTI in device
code
RTTI is not emitted for CUDA/HIP device code. As in OGCG, the RTTI
attribute returns a null pointer instead of a reference to the RTTI
descriptor, but dynamic_cast and typeid expected a GlobalViewAttr and
crashed on it.
Make the RTTI parameters of cir.dyn_cast_info optional. CIRGen leaves
them out when RTTI is not available, and the Itanium ABI lowering passes
null pointers to __dynamic_cast instead. The verifier requires either
both or neither, and since the attribute alias is named after the RTTI,
such an attribute is printed without one. typeid returns the null
pointer, like OGCG.
Signed-off-by: Steffen Holst Larsen <sholstla at amd.com>
---
.../include/clang/CIR/Dialect/IR/CIRAttrs.td | 10 +--
clang/lib/CIR/CodeGen/CIRGenExprCXX.cpp | 10 ++-
clang/lib/CIR/CodeGen/CIRGenItaniumCXXABI.cpp | 4 +-
clang/lib/CIR/Dialect/IR/CIRAttrs.cpp | 8 ++-
clang/lib/CIR/Dialect/IR/CIRDialect.cpp | 3 +
.../CIR/Dialect/Transforms/CXXABILowering.cpp | 16 +++--
.../TargetLowering/LowerItaniumCXXABI.cpp | 20 ++++--
clang/test/CIR/CodeGenCUDA/rtti-device.cu | 71 +++++++++++++++++++
clang/test/CIR/IR/dynamic-cast.cir | 12 ++++
clang/test/CIR/IR/invalid-dyn-cast.cir | 19 +++++
10 files changed, 152 insertions(+), 21 deletions(-)
create mode 100644 clang/test/CIR/CodeGenCUDA/rtti-device.cu
diff --git a/clang/include/clang/CIR/Dialect/IR/CIRAttrs.td b/clang/include/clang/CIR/Dialect/IR/CIRAttrs.td
index 358412b75fc05..95c9144b165d9 100644
--- a/clang/include/clang/CIR/Dialect/IR/CIRAttrs.td
+++ b/clang/include/clang/CIR/Dialect/IR/CIRAttrs.td
@@ -1201,7 +1201,9 @@ def CIR_DynamicCastInfoAttr : CIR_Attr<"DynamicCastInfo", "dyn_cast_info"> {
Provide ABI specific information about a dynamic cast operation.
The `src_rtti` and the `dest_rtti` parameters give the RTTI of the source
- record type and the destination record type, respectively.
+ record type and the destination record type, respectively. Both are absent
+ when RTTI is not emitted for the target (e.g. CUDA/HIP device code), in
+ which case null pointers are passed to the runtime function instead.
The `runtime_func` parameter gives the `__dynamic_cast` function which is
provided by the runtime. The `bad_cast_func` parameter gives the
@@ -1212,8 +1214,8 @@ def CIR_DynamicCastInfoAttr : CIR_Attr<"DynamicCastInfo", "dyn_cast_info"> {
}];
let parameters = (ins
- CIR_GlobalViewAttr:$src_rtti,
- CIR_GlobalViewAttr:$dest_rtti,
+ OptionalParameter<"cir::GlobalViewAttr">:$src_rtti,
+ OptionalParameter<"cir::GlobalViewAttr">:$dest_rtti,
"mlir::FlatSymbolRefAttr":$runtime_func,
"mlir::FlatSymbolRefAttr":$bad_cast_func,
CIR_IntAttr:$offset_hint
@@ -1226,7 +1228,7 @@ def CIR_DynamicCastInfoAttr : CIR_Attr<"DynamicCastInfo", "dyn_cast_info"> {
"mlir::FlatSymbolRefAttr":$runtime_func,
"mlir::FlatSymbolRefAttr":$bad_cast_func,
"IntAttr":$offset_hint), [{
- return $_get(src_rtti.getContext(), src_rtti, dest_rtti, runtime_func,
+ return $_get(runtime_func.getContext(), src_rtti, dest_rtti, runtime_func,
bad_cast_func, offset_hint);
}]>,
];
diff --git a/clang/lib/CIR/CodeGen/CIRGenExprCXX.cpp b/clang/lib/CIR/CodeGen/CIRGenExprCXX.cpp
index f8c2a5cce47e4..622653b76c418 100644
--- a/clang/lib/CIR/CodeGen/CIRGenExprCXX.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenExprCXX.cpp
@@ -1997,8 +1997,14 @@ mlir::Value CIRGenFunction::emitCXXTypeidExpr(const CXXTypeidExpr *e) {
// If the operand is already most derived object, no need to look up vtable.
}
- auto typeInfo =
- cast<cir::GlobalViewAttr>(cgm.getAddrOfRTTIDescriptor(loc, ty));
+ mlir::Attribute rtti = cgm.getAddrOfRTTIDescriptor(loc, ty);
+ // When RTTI is disabled for the target (e.g. CUDA/HIP device compilation),
+ // this is a null pointer instead of a reference to the descriptor.
+ auto typeInfo = dyn_cast<cir::GlobalViewAttr>(rtti);
+ if (!typeInfo)
+ return builder.createBitcast(
+ builder.getConstant(loc, cast<mlir::TypedAttr>(rtti)), resultType);
+
// `getAddrOfRTTIDescriptor` lies to us and always gives us a uint8ptr as its
// type, however we need the value of the actual global to call the
// get-global-op, so look it up here.
diff --git a/clang/lib/CIR/CodeGen/CIRGenItaniumCXXABI.cpp b/clang/lib/CIR/CodeGen/CIRGenItaniumCXXABI.cpp
index c7bf59871e691..9bc98a28b4afc 100644
--- a/clang/lib/CIR/CodeGen/CIRGenItaniumCXXABI.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenItaniumCXXABI.cpp
@@ -2293,9 +2293,9 @@ static cir::DynamicCastInfoAttr emitDynamicCastInfo(CIRGenFunction &cgf,
mlir::Location loc,
QualType srcRecordTy,
QualType destRecordTy) {
- auto srcRtti = mlir::cast<cir::GlobalViewAttr>(
+ auto srcRtti = mlir::dyn_cast<cir::GlobalViewAttr>(
cgf.cgm.getAddrOfRTTIDescriptor(loc, srcRecordTy.getUnqualifiedType()));
- auto destRtti = mlir::cast<cir::GlobalViewAttr>(
+ auto destRtti = mlir::dyn_cast<cir::GlobalViewAttr>(
cgf.cgm.getAddrOfRTTIDescriptor(loc, destRecordTy.getUnqualifiedType()));
cir::FuncOp runtimeFuncOp = getItaniumDynamicCastFn(cgf);
diff --git a/clang/lib/CIR/Dialect/IR/CIRAttrs.cpp b/clang/lib/CIR/Dialect/IR/CIRAttrs.cpp
index 2649ed1449739..705a445473ef4 100644
--- a/clang/lib/CIR/Dialect/IR/CIRAttrs.cpp
+++ b/clang/lib/CIR/Dialect/IR/CIRAttrs.cpp
@@ -998,10 +998,14 @@ LogicalResult DynamicCastInfoAttr::verify(
function_ref<InFlightDiagnostic()> emitError, cir::GlobalViewAttr srcRtti,
cir::GlobalViewAttr destRtti, mlir::FlatSymbolRefAttr runtimeFunc,
mlir::FlatSymbolRefAttr badCastFunc, cir::IntAttr offsetHint) {
- if (!isRttiPtr(srcRtti.getType()))
+ if (!srcRtti != !destRtti)
+ return emitError() << "srcRtti and destRtti must both be present or both "
+ "be absent";
+
+ if (srcRtti && !isRttiPtr(srcRtti.getType()))
return emitError() << "srcRtti must be an RTTI pointer";
- if (!isRttiPtr(destRtti.getType()))
+ if (destRtti && !isRttiPtr(destRtti.getType()))
return emitError() << "destRtti must be an RTTI pointer";
return success();
diff --git a/clang/lib/CIR/Dialect/IR/CIRDialect.cpp b/clang/lib/CIR/Dialect/IR/CIRDialect.cpp
index 4b9e0c7f0f544..2db94110a41bc 100644
--- a/clang/lib/CIR/Dialect/IR/CIRDialect.cpp
+++ b/clang/lib/CIR/Dialect/IR/CIRDialect.cpp
@@ -80,6 +80,9 @@ struct CIROpAsmDialectInterface : public OpAsmDialectInterface {
return AliasResult::FinalAlias;
}
if (auto dynCastInfoAttr = mlir::dyn_cast<cir::DynamicCastInfoAttr>(attr)) {
+ // The alias is named after the RTTI, so there is none without it.
+ if (!dynCastInfoAttr.getSrcRtti() || !dynCastInfoAttr.getDestRtti())
+ return AliasResult::NoAlias;
os << dynCastInfoAttr.getAlias();
return AliasResult::FinalAlias;
}
diff --git a/clang/lib/CIR/Dialect/Transforms/CXXABILowering.cpp b/clang/lib/CIR/Dialect/Transforms/CXXABILowering.cpp
index 9a0e40bd4f506..a534ee6676619 100644
--- a/clang/lib/CIR/Dialect/Transforms/CXXABILowering.cpp
+++ b/clang/lib/CIR/Dialect/Transforms/CXXABILowering.cpp
@@ -123,8 +123,10 @@ bool isCXXABIAttributeLegal(const mlir::TypeConverter &tc,
isCXXABIAttributeLegal(tc, tia.getData());
})
.Case<cir::DynamicCastInfoAttr>([&tc](cir::DynamicCastInfoAttr dcia) {
- return isCXXABIAttributeLegal(tc, dcia.getSrcRtti()) &&
- isCXXABIAttributeLegal(tc, dcia.getDestRtti()) &&
+ return (!dcia.getSrcRtti() ||
+ isCXXABIAttributeLegal(tc, dcia.getSrcRtti())) &&
+ (!dcia.getDestRtti() ||
+ isCXXABIAttributeLegal(tc, dcia.getDestRtti())) &&
isCXXABIAttributeLegal(tc, dcia.getRuntimeFunc()) &&
isCXXABIAttributeLegal(tc, dcia.getBadCastFunc());
})
@@ -230,11 +232,13 @@ mlir::Attribute rewriteAttribute(const mlir::TypeConverter &tc,
})
.Case<cir::DynamicCastInfoAttr>([&tc,
ctx](cir::DynamicCastInfoAttr dcia) {
+ auto rewriteRtti = [&](cir::GlobalViewAttr rtti) {
+ return rtti ? mlir::cast<cir::GlobalViewAttr>(
+ rewriteAttribute(tc, ctx, rtti))
+ : rtti;
+ };
return cir::DynamicCastInfoAttr::get(
- mlir::cast<cir::GlobalViewAttr>(
- rewriteAttribute(tc, ctx, dcia.getSrcRtti())),
- mlir::cast<cir::GlobalViewAttr>(
- rewriteAttribute(tc, ctx, dcia.getDestRtti())),
+ rewriteRtti(dcia.getSrcRtti()), rewriteRtti(dcia.getDestRtti()),
dcia.getRuntimeFunc(), dcia.getBadCastFunc(), dcia.getOffsetHint());
})
.Case<cir::ConstRecordAttr>([&tc, ctx](cir::ConstRecordAttr cra) {
diff --git a/clang/lib/CIR/Dialect/Transforms/TargetLowering/LowerItaniumCXXABI.cpp b/clang/lib/CIR/Dialect/Transforms/TargetLowering/LowerItaniumCXXABI.cpp
index 26dbe40c2c93f..e8c0a5bf1f36c 100644
--- a/clang/lib/CIR/Dialect/Transforms/TargetLowering/LowerItaniumCXXABI.cpp
+++ b/clang/lib/CIR/Dialect/Transforms/TargetLowering/LowerItaniumCXXABI.cpp
@@ -729,14 +729,24 @@ static mlir::Value buildDynamicCastAfterNullCheck(cir::DynamicCastOp op,
mlir::Value srcPtr = cir::CastOp::create(builder, loc, voidPtrTy,
cir::CastKind::bitcast, srcValue);
- mlir::Value srcRtti =
- cir::ConstantOp::create(builder, loc, castInfo.getSrcRtti());
- mlir::Value destRtti =
- cir::ConstantOp::create(builder, loc, castInfo.getDestRtti());
+ mlir::FlatSymbolRefAttr dynCastFuncRef = castInfo.getRuntimeFunc();
+ // Without RTTI for the target, the runtime function gets null pointers.
+ auto emitRtti = [&](cir::GlobalViewAttr rtti, unsigned argNo) -> mlir::Value {
+ if (rtti)
+ return cir::ConstantOp::create(builder, loc, rtti);
+ auto dynCastFunc = mlir::SymbolTable::lookupNearestSymbolFrom<cir::FuncOp>(
+ op, dynCastFuncRef);
+ auto rttiPtrTy = mlir::cast<cir::PointerType>(
+ dynCastFunc.getFunctionType().getInput(argNo));
+ return cir::ConstantOp::create(
+ builder, loc,
+ cir::ConstPtrAttr::get(rttiPtrTy, builder.getI64IntegerAttr(0)));
+ };
+ mlir::Value srcRtti = emitRtti(castInfo.getSrcRtti(), /*argNo=*/1);
+ mlir::Value destRtti = emitRtti(castInfo.getDestRtti(), /*argNo=*/2);
mlir::Value offsetHint =
cir::ConstantOp::create(builder, loc, castInfo.getOffsetHint());
- mlir::FlatSymbolRefAttr dynCastFuncRef = castInfo.getRuntimeFunc();
mlir::Value dynCastFuncArgs[4] = {srcPtr, srcRtti, destRtti, offsetHint};
mlir::Value castedPtr = cir::CallOp::create(builder, loc, dynCastFuncRef,
diff --git a/clang/test/CIR/CodeGenCUDA/rtti-device.cu b/clang/test/CIR/CodeGenCUDA/rtti-device.cu
new file mode 100644
index 0000000000000..79c2a24b1587a
--- /dev/null
+++ b/clang/test/CIR/CodeGenCUDA/rtti-device.cu
@@ -0,0 +1,71 @@
+#include "Inputs/cuda.h"
+
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -x cuda -fcuda-is-device -fclangir -emit-cir -mmlir --mlir-print-ir-before=cir-cxxabi-lowering %s -o /dev/null 2>&1 \
+// RUN: | FileCheck %s --check-prefix=CIR-BEFORE
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -x cuda -fcuda-is-device -fclangir -emit-cir %s -o - \
+// RUN: | FileCheck %s --check-prefix=CIR
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -x cuda -fcuda-is-device -fclangir -emit-llvm %s -o - \
+// RUN: | FileCheck %s --check-prefixes=LLVM,LLVMCIR
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -x cuda -fcuda-is-device -emit-llvm %s -o - \
+// RUN: | FileCheck %s --check-prefixes=LLVM,OGCG
+
+// RUN: %clang_cc1 -triple amdgpu9.00-amd-amdhsa -x hip -fcuda-is-device -fclangir -emit-cir -mmlir --mlir-print-ir-before=cir-cxxabi-lowering %s -o /dev/null 2>&1 \
+// RUN: | FileCheck %s --check-prefix=CIR-BEFORE
+// RUN: %clang_cc1 -triple amdgpu9.00-amd-amdhsa -x hip -fcuda-is-device -fclangir -emit-cir %s -o - \
+// RUN: | FileCheck %s --check-prefix=CIR
+// RUN: %clang_cc1 -triple amdgpu9.00-amd-amdhsa -x hip -fcuda-is-device -fclangir -emit-llvm %s -o - \
+// RUN: | FileCheck %s --check-prefixes=LLVM,LLVMCIR
+// RUN: %clang_cc1 -triple amdgpu9.00-amd-amdhsa -x hip -fcuda-is-device -emit-llvm %s -o - \
+// RUN: | FileCheck %s --check-prefixes=LLVM,OGCG
+
+namespace std {
+class type_info {
+public:
+ virtual ~type_info();
+};
+} // namespace std
+
+struct B {
+ __device__ virtual ~B();
+};
+struct D : B {};
+
+__device__ D *ptr_cast(B *b) { return dynamic_cast<D *>(b); }
+
+// CIR-BEFORE-LABEL: cir.func {{.*}} @_Z8ptr_castP1B(
+// CIR-BEFORE: cir.dyn_cast ptr %{{.+}} : !cir.ptr<!rec_B> -> !cir.ptr<!rec_D> #cir.dyn_cast_info<runtime_func = @__dynamic_cast, bad_cast_func = @__cxa_bad_cast, offset_hint = #cir.int<0> : !s64i>
+
+// CIR-LABEL: cir.func {{.*}} @_Z8ptr_castP1B(
+// CIR: %[[SRC_RTTI:.+]] = cir.const #cir.ptr<null> : !cir.ptr<!u8i>
+// CIR: %[[DEST_RTTI:.+]] = cir.const #cir.ptr<null> : !cir.ptr<!u8i>
+// CIR: cir.call @__dynamic_cast(%{{.+}}, %[[SRC_RTTI]], %[[DEST_RTTI]], %{{.+}})
+
+// LLVM-LABEL: define {{.*}} ptr @_Z8ptr_castP1B(
+// LLVM: call ptr @__dynamic_cast(ptr %{{.+}}, ptr {{.*}}null, ptr {{.*}}null, i64 0)
+
+__device__ D &ref_cast(B &b) { return dynamic_cast<D &>(b); }
+
+// CIR-BEFORE-LABEL: cir.func {{.*}} @_Z8ref_castR1B(
+// CIR-BEFORE: cir.dyn_cast ref %{{.+}} : !cir.ptr<!rec_B> -> !cir.ptr<!rec_D> #cir.dyn_cast_info<runtime_func = @__dynamic_cast, bad_cast_func = @__cxa_bad_cast, offset_hint = #cir.int<0> : !s64i>
+
+// CIR-LABEL: cir.func {{.*}} @_Z8ref_castR1B(
+// CIR: %[[SRC_RTTI:.+]] = cir.const #cir.ptr<null> : !cir.ptr<!u8i>
+// CIR: %[[DEST_RTTI:.+]] = cir.const #cir.ptr<null> : !cir.ptr<!u8i>
+// CIR: cir.call @__dynamic_cast(%{{.+}}, %[[SRC_RTTI]], %[[DEST_RTTI]], %{{.+}})
+// CIR: cir.call @__cxa_bad_cast()
+
+// LLVM-LABEL: define {{.*}} ptr @_Z8ref_castR1B(
+// LLVM: call ptr @__dynamic_cast(ptr %{{.+}}, ptr {{.*}}null, ptr {{.*}}null, i64 0)
+// LLVM: call void @__cxa_bad_cast()
+
+__device__ const std::type_info &static_typeid() { return typeid(D); }
+
+// CIR-LABEL: cir.func {{.*}} @_Z13static_typeidv(
+// CIR: %[[RTTI:.+]] = cir.const #cir.ptr<null> : !cir.ptr<!u8i>
+// CIR: cir.cast bitcast %[[RTTI]] : !cir.ptr<!u8i> -> !cir.ptr<!rec_std3A3Atype_info>
+
+// LLVM-LABEL: define {{.*}} ptr @_Z13static_typeidv(
+// LLVMCIR: store ptr null, ptr {{(addrspace\(5\) )?}}[[RET:%.+]], align 8
+// LLVMCIR: [[RES:%.+]] = load ptr, ptr {{(addrspace\(5\) )?}}[[RET]], align 8
+// LLVMCIR: ret ptr [[RES]]
+// OGCG: ret ptr {{null|addrspacecast \(ptr addrspace\(1\) null to ptr\)}}
diff --git a/clang/test/CIR/IR/dynamic-cast.cir b/clang/test/CIR/IR/dynamic-cast.cir
index 6930db82ca0da..aa2b73a01c39e 100644
--- a/clang/test/CIR/IR/dynamic-cast.cir
+++ b/clang/test/CIR/IR/dynamic-cast.cir
@@ -37,6 +37,18 @@ module {
// CHECK: cir.return %0 : !cir.ptr<!rec_Derived>
// CHECK: }
+ // Without RTTI for the target (e.g. CUDA/HIP device code) the RTTI is
+ // absent, and so is the alias.
+ cir.func @test_no_rtti_cast(%arg0: !cir.ptr<!rec_Base>) -> !cir.ptr<!rec_Derived> {
+ %0 = cir.dyn_cast ptr %arg0 : !cir.ptr<!rec_Base> -> !cir.ptr<!rec_Derived> #cir.dyn_cast_info<runtime_func = @__dynamic_cast, bad_cast_func = @__cxa_bad_cast, offset_hint = #cir.int<0> : !s64i>
+ cir.return %0 : !cir.ptr<!rec_Derived>
+ }
+
+ // CHECK: cir.func @test_no_rtti_cast(%arg0: !cir.ptr<!rec_Base>) -> !cir.ptr<!rec_Derived> {
+ // CHECK: %0 = cir.dyn_cast ptr %arg0 : !cir.ptr<!rec_Base> -> !cir.ptr<!rec_Derived> #cir.dyn_cast_info<runtime_func = @__dynamic_cast, bad_cast_func = @__cxa_bad_cast, offset_hint = #cir.int<0> : !s64i>
+ // CHECK: cir.return %0 : !cir.ptr<!rec_Derived>
+ // CHECK: }
+
cir.func dso_local @test_cast_to_void(%arg0: !cir.ptr<!rec_Base>) -> !cir.ptr<!void> {
%0 = cir.dyn_cast ptr %arg0 : !cir.ptr<!rec_Base> -> !cir.ptr<!void>
cir.return %0 : !cir.ptr<!void>
diff --git a/clang/test/CIR/IR/invalid-dyn-cast.cir b/clang/test/CIR/IR/invalid-dyn-cast.cir
index c5c7df784d4fc..346ca78d888fd 100644
--- a/clang/test/CIR/IR/invalid-dyn-cast.cir
+++ b/clang/test/CIR/IR/invalid-dyn-cast.cir
@@ -41,3 +41,22 @@ module {
%0 = cir.dyn_cast ptr %arg0 : !cir.ptr<!Base> -> !cir.ptr<!Derived> #cir.dyn_cast_info<src_rtti = #cir.global_view<@_ZTI4Base> : !cir.ptr<!u8i>, dest_rtti = #cir.global_view<@_ZTI7Derived> : !cir.ptr<!u32i>, runtime_func = @__dynamic_cast, bad_cast_func = @__cxa_bad_cast, offset_hint = #cir.int<0> : !s64i>
}
}
+
+// -----
+
+!s64i = !cir.int<s, 64>
+!u8i = !cir.int<u, 8>
+!void = !cir.void
+
+!Base = !cir.struct<"Base" {data !cir.ptr<!cir.ptr<!cir.func<() -> !cir.int<u, 32>>>>}>
+!Derived = !cir.struct<"Derived" {data !cir.struct<"Base" {data !cir.ptr<!cir.ptr<!cir.func<() -> !cir.int<u, 32>>>>}>}>
+
+module {
+ cir.global "private" constant external @_ZTI4Base : !cir.ptr<!u8i>
+ cir.func private @__dynamic_cast(!cir.ptr<!void>, !cir.ptr<!u8i>, !cir.ptr<!u8i>, !s64i) -> !cir.ptr<!void>
+ cir.func private @__cxa_bad_cast()
+ cir.func @test(%arg0 : !cir.ptr<!Base>) {
+ // expected-error at +1 {{srcRtti and destRtti must both be present or both be absent}}
+ %0 = cir.dyn_cast ptr %arg0 : !cir.ptr<!Base> -> !cir.ptr<!Derived> #cir.dyn_cast_info<src_rtti = #cir.global_view<@_ZTI4Base> : !cir.ptr<!u8i>, runtime_func = @__dynamic_cast, bad_cast_func = @__cxa_bad_cast, offset_hint = #cir.int<0> : !s64i>
+ }
+}
More information about the cfe-commits
mailing list