• Home
  • Line#
  • Scopes#
  • Navigate#
  • Raw
  • Download
1 // CUDA struct types with interesting initialization properties.
2 // Keep in sync with ../SemaCUDA/Inputs/cuda-initializers.h.
3 
4 // Base classes with different initializer variants.
5 
6 // trivial constructor -- allowed
7 struct T {
8   int t;
9 };
10 
11 // empty constructor
12 struct EC {
13   int ec;
ECEC14   __device__ EC() {}     // -- allowed
ECEC15   __device__ EC(int) {}  // -- not allowed
16 };
17 
18 // empty destructor
19 struct ED {
~EDED20   __device__ ~ED() {}     // -- allowed
21 };
22 
23 struct ECD {
ECDECD24   __device__ ECD() {}     // -- allowed
~ECDECD25   __device__ ~ECD() {}    // -- allowed
26 };
27 
28 // empty templated constructor -- allowed with no arguments
29 struct ETC {
ETCETC30   template <typename... T> __device__ ETC(T...) {}
31 };
32 
33 // undefined constructor -- not allowed
34 struct UC {
35   int uc;
36   __device__ UC();
37 };
38 
39 // undefined destructor -- not allowed
40 struct UD {
41   int ud;
42   __device__ ~UD();
43 };
44 
45 // empty constructor w/ initializer list -- not allowed
46 struct ECI {
47   int eci;
ECIECI48   __device__ ECI() : eci(1) {}
49 };
50 
51 // non-empty constructor -- not allowed
52 struct NEC {
53   int nec;
NECNEC54   __device__ NEC() { nec = 1; }
55 };
56 
57 // non-empty destructor -- not allowed
58 struct NED {
59   int ned;
~NEDNED60   __device__ ~NED() { ned = 1; }
61 };
62 
63 // no-constructor,  virtual method -- not allowed
64 struct NCV {
65   int ncv;
vmNCV66   __device__ virtual void vm() {}
67 };
68 
69 // virtual destructor -- not allowed.
70 struct VD {
~VDVD71   __device__ virtual ~VD() {}
72 };
73 
74 // dynamic in-class field initializer -- not allowed
75 __device__ int f();
76 struct NCF {
77   int ncf = f();
78 };
79 
80 // static in-class field initializer.  NVCC does not allow it, but
81 // clang generates static initializer for this, so we'll accept it.
82 // We still can't use it on __shared__ vars as they don't allow *any*
83 // initializers.
84 struct NCFS {
85   int ncfs = 3;
86 };
87 
88 // undefined templated constructor -- not allowed
89 struct UTC {
90   template <typename... T> __device__ UTC(T...);
91 };
92 
93 // non-empty templated constructor -- not allowed
94 struct NETC {
95   int netc;
NETCNETC96   template <typename... T> __device__ NETC(T...) { netc = 1; }
97 };
98 
99 // Regular base class -- allowed
100 struct T_B_T : T {};
101 
102 // Incapsulated object of allowed class -- allowed
103 struct T_F_T {
104   T t;
105 };
106 
107 // array of allowed objects -- allowed
108 struct T_FA_T {
109   T t[2];
110 };
111 
112 
113 // Calling empty base class initializer is OK
114 struct EC_I_EC : EC {
EC_I_ECEC_I_EC115   __device__ EC_I_EC() : EC() {}
116 };
117 
118 // .. though passing arguments is not allowed.
119 struct EC_I_EC1 : EC {
EC_I_EC1EC_I_EC1120   __device__ EC_I_EC1() : EC(1) {}
121 };
122 
123 // Virtual base class -- not allowed
124 struct T_V_T : virtual T {};
125 
126 // Inherited from or incapsulated class with non-empty constructor --
127 // not allowed
128 struct T_B_NEC : NEC {};
129 struct T_F_NEC {
130   NEC nec;
131 };
132 struct T_FA_NEC {
133   NEC nec[2];
134 };
135 
136 
137 // Inherited from or incapsulated class with non-empty desstructor --
138 // not allowed
139 struct T_B_NED : NED {};
140 struct T_F_NED {
141   NED ned;
142 };
143 struct T_FA_NED {
144   NED ned[2];
145 };
146