69 lines · plain
1// REQUIRES: x86-registered-target2// REQUIRES: nvptx-registered-target3 4// Make sure we don't emit vtables for classes with methods that have5// inappropriate target attributes. Currently it's mostly needed in6// order to avoid emitting vtables for host-only classes on device7// 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-BOTH11// 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-BOTH13 14#include "Inputs/cuda.h"15 16struct H {17 virtual void method();18};19//CHECK-HOST: @_ZTV1H =20//CHECK-HOST-SAME: @_ZN1H6methodEv21//CHECK-DEVICE-NOT: @_ZTV1H =22//CHECK-DEVICE-NOT: @_ZTVN10__cxxabiv117__class_type_infoE23//CHECK-DEVICE-NOT: @_ZTS1H24//CHECK-DEVICE-NOT: @_ZTI1H25struct D {26 __device__ virtual void method();27};28 29//CHECK-DEVICE: @_ZTV1D30//CHECK-DEVICE-SAME: @_ZN1D6methodEv31//CHECK-HOST-NOT: @_ZTV1D32//CHECK-DEVICE-NOT: @_ZTVN10__cxxabiv117__class_type_infoE33//CHECK-DEVICE-NOT: @_ZTS1D34//CHECK-DEVICE-NOT: @_ZTI1D35// This is the case with mixed host and device virtual methods. It's36// impossible to emit a valid vtable in that case because only host or37// only device methods would be available during host or device38// compilation. At the moment Clang (and NVCC) emit NULL pointers for39// unavailable methods,40struct HD {41 virtual void h_method();42 __device__ virtual void d_method();43};44// CHECK-BOTH: @_ZTV2HD45// CHECK-DEVICE-NOT: @_ZN2HD8h_methodEv46// CHECK-DEVICE-SAME: null47// CHECK-DEVICE-SAME: @_ZN2HD8d_methodEv48// CHECK-HOST-SAME: @_ZN2HD8h_methodEv49// CHECK-HOST-NOT: @_ZN2HD8d_methodEv50// CHECK-HOST-SAME: null51// CHECK-BOTH-SAME: ]52// CHECK-DEVICE-NOT: @_ZTVN10__cxxabiv117__class_type_infoE53// CHECK-DEVICE-NOT: @_ZTS2HD54// CHECK-DEVICE-NOT: @_ZTI2HD55 56void H::method() {}57//CHECK-HOST: define{{.*}} void @_ZN1H6methodEv58 59void __device__ D::method() {}60//CHECK-DEVICE: define{{.*}} void @_ZN1D6methodEv61 62void __device__ HD::d_method() {}63// CHECK-DEVICE: define{{.*}} void @_ZN2HD8d_methodEv64// CHECK-HOST-NOT: define{{.*}} void @_ZN2HD8d_methodEv65void HD::h_method() {}66// CHECK-HOST: define{{.*}} void @_ZN2HD8h_methodEv67// CHECK-DEVICE-NOT: define{{.*}} void @_ZN2HD8h_methodEv68 69