1 // REQUIRES: x86-registered-target
2 // REQUIRES: nvptx-registered-target
3
4 // Make sure we handle target overloads correctly. Most of this is checked in
5 // sema, but special functions like constructors and destructors are here.
6 //
7 // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm -o - %s \
8 // RUN: | FileCheck -check-prefix=CHECK-BOTH -check-prefix=CHECK-HOST %s
9 // RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -emit-llvm -o - %s \
10 // RUN: | FileCheck -check-prefix=CHECK-BOTH -check-prefix=CHECK-DEVICE %s
11
12 #include "Inputs/cuda.h"
13
14 // Check constructors/destructors for D/H functions
15 int x;
16 struct s_cd_dh {
s_cd_dhs_cd_dh17 __host__ s_cd_dh() { x = 11; }
s_cd_dhs_cd_dh18 __device__ s_cd_dh() { x = 12; }
~s_cd_dhs_cd_dh19 __host__ ~s_cd_dh() { x = 21; }
~s_cd_dhs_cd_dh20 __device__ ~s_cd_dh() { x = 22; }
21 };
22
23 struct s_cd_hd {
s_cd_hds_cd_hd24 __host__ __device__ s_cd_hd() { x = 31; }
~s_cd_hds_cd_hd25 __host__ __device__ ~s_cd_hd() { x = 32; }
26 };
27
28 // CHECK-BOTH: define void @_Z7wrapperv
29 #if defined(__CUDA_ARCH__)
30 __device__
31 #else
32 __host__
33 #endif
wrapper()34 void wrapper() {
35 s_cd_dh scddh;
36 // CHECK-BOTH: call void @_ZN7s_cd_dhC1Ev(
37 s_cd_hd scdhd;
38 // CHECK-BOTH: call void @_ZN7s_cd_hdC1Ev
39
40 // CHECK-BOTH: call void @_ZN7s_cd_hdD1Ev(
41 // CHECK-BOTH: call void @_ZN7s_cd_dhD1Ev(
42 }
43 // CHECK-BOTH: ret void
44
45 // Now it's time to check what's been generated for the methods we used.
46
47 // CHECK-BOTH: define linkonce_odr void @_ZN7s_cd_dhC2Ev(
48 // CHECK-HOST: store i32 11,
49 // CHECK-DEVICE: store i32 12,
50 // CHECK-BOTH: ret void
51
52 // CHECK-BOTH: define linkonce_odr void @_ZN7s_cd_hdC2Ev(
53 // CHECK-BOTH: store i32 31,
54 // CHECK-BOTH: ret void
55
56 // CHECK-BOTH: define linkonce_odr void @_ZN7s_cd_hdD2Ev(
57 // CHECK-BOTH: store i32 32,
58 // CHECK-BOTH: ret void
59
60 // CHECK-BOTH: define linkonce_odr void @_ZN7s_cd_dhD2Ev(
61 // CHECK-HOST: store i32 21,
62 // CHECK-DEVICE: store i32 22,
63 // CHECK-BOTH: ret void
64