[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:

I am not sure I fully agree that the link error isn't fairly obvious, given it rejects the specific operations by name. I do agree that a diagnostic would be clearer and a quick experiment shows that it also matches NVCC. I can propose that solution separately, probably in Sema. That should reject it in both. It is a behavioral change and there might be a reason the current solution is done historically, so I can't guarantee that it will gain any traction.

That just leaves the question of whether we still want to match OGCG here. It would be the safer option, as it means we will have a solution to RTTI on device that matches OGCG and won't have to return here if the rejection diagnostic solution isn't accepted. On the other hand, it probably won't have any benefit if it is accepted.

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