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
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
|
// RUN: %clang_cc1 -x hip %s --hipstdpar -triple amdgcn-amd-amdhsa --std=c++17 \
// RUN: -fcuda-is-device -emit-llvm -o /dev/null -verify
// Note: These would happen implicitly, within the implementation of the
// accelerator specific algorithm library, and not from user code.
// Calls from the accelerator side to implicitly host (i.e. unannotated)
// functions are fine.
// expected-no-diagnostics
#define __device__ __attribute__((device))
#define __global__ __attribute__((global))
extern "C" void host_fn() {}
struct Dummy {};
struct S {
S() {}
~S() { host_fn(); }
int x;
};
struct T {
__device__ void hd() { host_fn(); }
__device__ void hd3();
void h() {}
void operator+();
void operator-(const T&) {}
operator Dummy() { return Dummy(); }
};
__device__ void T::hd3() { host_fn(); }
template <typename T> __device__ void hd2() { host_fn(); }
__global__ void kernel() { hd2<int>(); }
__device__ void hd() { host_fn(); }
template <typename T> __device__ void hd3() { host_fn(); }
__device__ void device_fn() { hd3<int>(); }
__device__ void local_var() {
S s;
}
__device__ void explicit_destructor(S *s) {
s->~S();
}
__device__ void hd_member_fn() {
T t;
t.hd();
}
__device__ void h_member_fn() {
T t;
t.h();
}
__device__ void unaryOp() {
T t;
(void) +t;
}
__device__ void binaryOp() {
T t;
(void) (t - t);
}
__device__ void implicitConversion() {
T t;
Dummy d = t;
}
template <typename T>
struct TmplStruct {
template <typename U> __device__ void fn() {}
};
template <>
template <>
__device__ void TmplStruct<int>::fn<int>() { host_fn(); }
__device__ void double_specialization() { TmplStruct<int>().fn<int>(); }
|
