1 // REQUIRES: x86-registered-target
2 // REQUIRES: nvptx-registered-target
3 
4 // Make sure we don't emit vtables for classes with methods that have
5 // inappropriate target attributes. Currently it's mostly needed in
6 // order to avoid emitting vtables for host-only classes on device
7 // side where we can't codegen them.
8 
9 // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm -o - %s \
10 // RUN:     | FileCheck %s -check-prefix=CHECK-HOST -check-prefix=CHECK-BOTH
11 // RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -emit-llvm -o - %s \
12 // RUN:     | FileCheck %s -check-prefix=CHECK-DEVICE -check-prefix=CHECK-BOTH
13 
14 #include "Inputs/cuda.h"
15 
16 struct H  {
17   virtual void method();
18 };
19 //CHECK-HOST: @_ZTV1H =
20 //CHECK-HOST-SAME: @_ZN1H6methodEv
21 //CHECK-DEVICE-NOT: @_ZTV1H =
22 //CHECK-DEVICE-NOT: @_ZTVN10__cxxabiv117__class_type_infoE
23 //CHECK-DEVICE-NOT: @_ZTS1H
24 //CHECK-DEVICE-NOT: @_ZTI1H
25 struct D  {
26    __device__ virtual void method();
27 };
28 
29 //CHECK-DEVICE: @_ZTV1D
30 //CHECK-DEVICE-SAME: @_ZN1D6methodEv
31 //CHECK-HOST-NOT: @_ZTV1D
32 //CHECK-DEVICE-NOT: @_ZTVN10__cxxabiv117__class_type_infoE
33 //CHECK-DEVICE-NOT: @_ZTS1D
34 //CHECK-DEVICE-NOT: @_ZTI1D
35 // This is the case with mixed host and device virtual methods.  It's
36 // impossible to emit a valid vtable in that case because only host or
37 // only device methods would be available during host or device
38 // compilation. At the moment Clang (and NVCC) emit NULL pointers for
39 // unavailable methods,
40 struct HD  {
41   virtual void h_method();
42   __device__ virtual void d_method();
43 };
44 // CHECK-BOTH: @_ZTV2HD
45 // CHECK-DEVICE-NOT: @_ZN2HD8h_methodEv
46 // CHECK-DEVICE-SAME: null
47 // CHECK-DEVICE-SAME: @_ZN2HD8d_methodEv
48 // CHECK-HOST-SAME: @_ZN2HD8h_methodEv
49 // CHECK-HOST-NOT: @_ZN2HD8d_methodEv
50 // CHECK-HOST-SAME: null
51 // CHECK-BOTH-SAME: ]
52 // CHECK-DEVICE-NOT: @_ZTVN10__cxxabiv117__class_type_infoE
53 // CHECK-DEVICE-NOT: @_ZTS2HD
54 // CHECK-DEVICE-NOT: @_ZTI2HD
55 
method()56 void H::method() {}
57 //CHECK-HOST: define void @_ZN1H6methodEv
58 
method()59 void __device__ D::method() {}
60 //CHECK-DEVICE: define void @_ZN1D6methodEv
61 
d_method()62 void __device__ HD::d_method() {}
63 // CHECK-DEVICE: define void @_ZN2HD8d_methodEv
64 // CHECK-HOST-NOT: define void @_ZN2HD8d_methodEv
h_method()65 void HD::h_method() {}
66 // CHECK-HOST: define void @_ZN2HD8h_methodEv
67 // CHECK-DEVICE-NOT: define void @_ZN2HD8h_methodEv
68 
69