================
@@ -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
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
