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 56 void H::method() {} 57 //CHECK-HOST: define void @_ZN1H6methodEv 58 59 void __device__ D::method() {} 60 //CHECK-DEVICE: define void @_ZN1D6methodEv 61 62 void __device__ HD::d_method() {} 63 // CHECK-DEVICE: define void @_ZN2HD8d_methodEv 64 // CHECK-HOST-NOT: define void @_ZN2HD8d_methodEv 65 void HD::h_method() {} 66 // CHECK-HOST: define void @_ZN2HD8h_methodEv 67 // CHECK-DEVICE-NOT: define void @_ZN2HD8h_methodEv 68 69