[clang] [CIR] Handle dynamic_cast and typeid without RTTI in device code (PR #228031)

[Steffen Larsen via cfe-commits]

================
@@ -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); }
----------------
steffenlarsen wrote:

It is indeed not supported in device code and even though we let it through here it will fail to link in device code. That said, there are two reasons why we would want this:

1. Currently there are cases where CIR places certain host-only vtable entries in device code as well. If the entries contain RTTI the device compilation fails.
2. The proposed solution matches OGCG and arguably the linking error is clearer than a compiler crash.

Even if we fix 1. (which I am looking into), I think 2. still stands. We could consider giving an error if we find RTTI in the device code, but I think the linking error is clear enough and we don't risk other unforeseen issues from the deviance from OGCG behavior.

https://github.com/llvm/llvm-project/pull/228031