================
@@ -0,0 +1,100 @@
+// RUN: %clang_cc1 -fcuda-is-device -fsyntax-only -verify=expected,dev %s
+// RUN: %clang_cc1 -fsyntax-only -verify %s
+// RUN: %clang_cc1 -x hip -fcuda-is-device -fsyntax-only -verify=expected,dev
%s
+// RUN: %clang_cc1 -x hip -fsyntax-only -verify %s
+
+#include "Inputs/cuda.h"
+
+namespace std {
+class type_info {};
+} // namespace std
+
+struct B {
+ __host__ __device__ virtual ~B() {}
+};
+struct D : B {};
+
+void host(B *b) {
+ (void)dynamic_cast<D *>(b);
+ (void)typeid(*b);
+}
+
+__device__ void device(B *b, D *d) {
+ (void)dynamic_cast<D *>(b);
+ // expected-error@-1 {{cannot use 'dynamic_cast' in __device__ function as
RTTI is not available in device code}}
+ (void)dynamic_cast<D &>(*b);
+ // expected-error@-1 {{cannot use 'dynamic_cast' in __device__ function as
RTTI is not available in device code}}
+ (void)typeid(D);
+ // expected-error@-1 {{cannot use 'typeid' in __device__ function as RTTI is
not available in device code}}
+ (void)typeid(*b);
+ // expected-error@-1 {{cannot use 'typeid' in __device__ function as RTTI is
not available in device code}}
+ (void)sizeof(typeid(int));
+ // expected-error@-1 {{cannot use 'typeid' in __device__ function as RTTI is
not available in device code}}
+
+ // As with -fno-rtti, these don't use RTTI and are allowed.
+ (void)dynamic_cast<void *>(b);
+ (void)dynamic_cast<B *>(d);
+}
+
+__global__ void kernel(B *b) {
+ (void)typeid(*b);
+ // expected-error@-1 {{cannot use 'typeid' in __global__ function as RTTI is
not available in device code}}
+}
+
+// Check that it's an error to use RTTI from a __host__ __device__ function if
+// and only if it's codegen'ed for device.
+
+__host__ __device__ void hd1(B *b) {
+ (void)dynamic_cast<D *>(b);
+ // dev-error@-1 {{cannot use 'dynamic_cast' in __host__ __device__ function
as RTTI is not available in device code}}
+}
+
+// No error, never instantiated on device.
----------------
steffenlarsen wrote:
This seems to be a wider gap than just RTTI, so I made it a separate PR:
https://github.com/llvm/llvm-project/pull/229008
Depending on which gets merged first, the other should add the RTTI cases.
https://github.com/llvm/llvm-project/pull/228469
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
