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

Reply via email to