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