| 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 | |
|---|