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 { |
17 | __host__ s_cd_dh() { x = 11; } |
18 | __device__ s_cd_dh() { x = 12; } |
19 | }; |
20 | |
21 | struct s_cd_hd { |
22 | __host__ __device__ s_cd_hd() { x = 31; } |
23 | __host__ __device__ ~s_cd_hd() { x = 32; } |
24 | }; |
25 | |
26 | // CHECK-BOTH: define void @_Z7wrapperv |
27 | #if defined(__CUDA_ARCH__) |
28 | __device__ |
29 | #else |
30 | __host__ |
31 | #endif |
32 | void wrapper() { |
33 | s_cd_dh scddh; |
34 | // CHECK-BOTH: call void @_ZN7s_cd_dhC1Ev( |
35 | s_cd_hd scdhd; |
36 | // CHECK-BOTH: call void @_ZN7s_cd_hdC1Ev |
37 | |
38 | // CHECK-BOTH: call void @_ZN7s_cd_hdD1Ev( |
39 | } |
40 | // CHECK-BOTH: ret void |
41 | |
42 | // Now it's time to check what's been generated for the methods we used. |
43 | |
44 | // CHECK-BOTH: define linkonce_odr void @_ZN7s_cd_dhC2Ev( |
45 | // CHECK-HOST: store i32 11, |
46 | // CHECK-DEVICE: store i32 12, |
47 | // CHECK-BOTH: ret void |
48 | |
49 | // CHECK-BOTH: define linkonce_odr void @_ZN7s_cd_hdC2Ev( |
50 | // CHECK-BOTH: store i32 31, |
51 | // CHECK-BOTH: ret void |
52 | |
53 | // CHECK-BOTH: define linkonce_odr void @_ZN7s_cd_hdD2Ev( |
54 | // CHECK-BOTH: store i32 32, |
55 | // CHECK-BOTH: ret void |
56 | |