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
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
| // CUDA struct types with interesting initialization properties.
// Keep in sync with ../SemaCUDA/Inputs/cuda-initializers.h.
// Base classes with different initializer variants.
// trivial constructor -- allowed
struct T {
int t;
};
// empty constructor
struct EC {
int ec;
__device__ EC() {} // -- allowed
__device__ EC(int) {} // -- not allowed
};
// empty destructor
struct ED {
__device__ ~ED() {} // -- allowed
};
struct ECD {
__device__ ECD() {} // -- allowed
__device__ ~ECD() {} // -- allowed
};
// empty templated constructor -- allowed with no arguments
struct ETC {
template <typename... T> __device__ ETC(T...) {}
};
// undefined constructor -- not allowed
struct UC {
int uc;
__device__ UC();
};
// undefined destructor -- not allowed
struct UD {
int ud;
__device__ ~UD();
};
// empty constructor w/ initializer list -- not allowed
struct ECI {
int eci;
__device__ ECI() : eci(1) {}
};
// non-empty constructor -- not allowed
struct NEC {
int nec;
__device__ NEC() { nec = 1; }
};
// non-empty destructor -- not allowed
struct NED {
int ned;
__device__ ~NED() { ned = 1; }
};
// no-constructor, virtual method -- not allowed
struct NCV {
int ncv;
__device__ virtual void vm() {}
};
// virtual destructor -- not allowed.
struct VD {
__device__ virtual ~VD() {}
};
// dynamic in-class field initializer -- not allowed
__device__ int f();
struct NCF {
int ncf = f();
};
// static in-class field initializer. NVCC does not allow it, but
// clang generates static initializer for this, so we'll accept it.
// We still can't use it on __shared__ vars as they don't allow *any*
// initializers.
struct NCFS {
int ncfs = 3;
};
// undefined templated constructor -- not allowed
struct UTC {
template <typename... T> __device__ UTC(T...);
};
// non-empty templated constructor -- not allowed
struct NETC {
int netc;
template <typename... T> __device__ NETC(T...) { netc = 1; }
};
// Regular base class -- allowed
struct T_B_T : T {};
// Incapsulated object of allowed class -- allowed
struct T_F_T {
T t;
};
// array of allowed objects -- allowed
struct T_FA_T {
T t[2];
};
// Calling empty base class initializer is OK
struct EC_I_EC : EC {
__device__ EC_I_EC() : EC() {}
};
// .. though passing arguments is not allowed.
struct EC_I_EC1 : EC {
__device__ EC_I_EC1() : EC(1) {}
};
// Virtual base class -- not allowed
struct T_V_T : virtual T {};
// Inherited from or incapsulated class with non-empty constructor --
// not allowed
struct T_B_NEC : NEC {};
struct T_F_NEC {
NEC nec;
};
struct T_FA_NEC {
NEC nec[2];
};
// Inherited from or incapsulated class with non-empty desstructor --
// not allowed
struct T_B_NED : NED {};
struct T_F_NED {
NED ned;
};
struct T_FA_NED {
NED ned[2];
};
|