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