Avoid depending on test inputes that aren't in Inputs

Some people have weird CI systems that run each test subdirectory
independently without access to other parallel trees.

Unfortunately, this means we have to suffer some duplication until Art
can sort out how to share these types.

git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@270164 91177308-0d34-0410-b5e6-96231b3b80d8

2 files changed, 157 insertions, 133 deletions

diff --git a/test/CodeGenCUDA/Inputs/cuda-initializers.h b/test/CodeGenCUDA/Inputs/cuda-initializers.h
new file mode 100644
index 0000000000..186b160276
--- /dev/null
+++ b/test/CodeGenCUDA/Inputs/cuda-initializers.h
@@ -0,0 +1,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];
+};
diff --git a/test/CodeGenCUDA/device-var-init.cu b/test/CodeGenCUDA/device-var-init.cu
index fa24dd95f9..6f2d929413 100644
--- a/test/CodeGenCUDA/device-var-init.cu
+++ b/test/CodeGenCUDA/device-var-init.cu
@@ -10,100 +10,8 @@
 #include "Inputs/cuda.h"
 #endif
 
-// 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; }
-};
+// Use the types we share with Sema tests.
+#include "Inputs/cuda-initializers.h"
 
 __device__ int d_v;
 // CHECK: @d_v = addrspace(1) externally_initialized global i32 0,
@@ -115,6 +23,7 @@ __constant__ int c_v;
 __device__ int d_v_i = 1;
 // CHECK: @d_v_i = addrspace(1) externally_initialized global i32 1,
 
+// trivial constructor -- allowed
 __device__ T d_t;
 // CHECK: @d_t = addrspace(1) externally_initialized global %struct.T zeroinitializer
 __shared__ T s_t;
@@ -127,6 +36,7 @@ __device__ T d_t_i = {2};
 __constant__ T c_t_i = {2};
 // CHECK: @c_t_i = addrspace(4) externally_initialized global %struct.T { i32 2 },
 
+// empty constructor
 __device__ EC d_ec;
 // CHECK: @d_ec = addrspace(1) externally_initialized global %struct.EC zeroinitializer,
 __shared__ EC s_ec;
@@ -134,6 +44,7 @@ __shared__ EC s_ec;
 __constant__ EC c_ec;
 // CHECK: @c_ec = addrspace(4) externally_initialized global %struct.EC zeroinitializer,
 
+// empty destructor
 __device__ ED d_ed;
 // CHECK: @d_ed = addrspace(1) externally_initialized global %struct.ED zeroinitializer,
 __shared__ ED s_ed;
@@ -148,6 +59,7 @@ __shared__ ECD s_ecd;
 __constant__ ECD c_ecd;
 // CHECK: @c_ecd = addrspace(4) externally_initialized global %struct.ECD zeroinitializer,
 
+// empty templated constructor -- allowed with no arguments
 __device__ ETC d_etc;
 // CHECK: @d_etc = addrspace(1) externally_initialized global %struct.ETC zeroinitializer,
 __shared__ ETC s_etc;
@@ -161,7 +73,6 @@ __constant__ NCFS c_ncfs;
 // CHECK: @c_ncfs = addrspace(4) externally_initialized global %struct.NCFS { i32 3 }
 
 // Regular base class -- allowed
-struct T_B_T : T {};
 __device__ T_B_T d_t_b_t;
 // CHECK: @d_t_b_t = addrspace(1) externally_initialized global %struct.T_B_T zeroinitializer,
 __shared__ T_B_T s_t_b_t;
@@ -170,9 +81,6 @@ __constant__ T_B_T c_t_b_t;
 // CHECK: @c_t_b_t = addrspace(4) externally_initialized global %struct.T_B_T zeroinitializer,
 
 // Incapsulated object of allowed class -- allowed
-struct T_F_T {
-  T t;
-};
 __device__ T_F_T d_t_f_t;
 // CHECK: @d_t_f_t = addrspace(1) externally_initialized global %struct.T_F_T zeroinitializer,
 __shared__ T_F_T s_t_f_t;
@@ -181,9 +89,6 @@ __constant__ T_F_T c_t_f_t;
 // CHECK: @c_t_f_t = addrspace(4) externally_initialized global %struct.T_F_T zeroinitializer,
 
 // array of allowed objects -- allowed
-struct T_FA_T {
-  T t[2];
-};
 __device__ T_FA_T d_t_fa_t;
 // CHECK: @d_t_fa_t = addrspace(1) externally_initialized global %struct.T_FA_T zeroinitializer,
 __shared__ T_FA_T s_t_fa_t;
@@ -193,9 +98,6 @@ __constant__ T_FA_T c_t_fa_t;
 
 
 // Calling empty base class initializer is OK
-struct EC_I_EC : EC {
-  __device__ EC_I_EC() : EC() {}
-};
 __device__ EC_I_EC d_ec_i_ec;
 // CHECK: @d_ec_i_ec = addrspace(1) externally_initialized global %struct.EC_I_EC zeroinitializer,
 __shared__ EC_I_EC s_ec_i_ec;
@@ -203,35 +105,6 @@ __shared__ EC_I_EC s_ec_i_ec;
 __constant__ EC_I_EC c_ec_i_ec;
 // CHECK: @c_ec_i_ec = addrspace(4) externally_initialized global %struct.EC_I_EC zeroinitializer,
 
-// .. 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];
-};
-
 // We should not emit global initializers for device-side variables.
 // CHECK-NOT: @__cxx_global_var_init
 
@@ -249,14 +122,20 @@ __device__ void df() {
   ETC etc;
   // CHECK:   call void @_ZN3ETCC1IJEEEDpT_(%struct.ETC* %etc)
   UC uc;
+  // undefined constructor -- not allowed
   // CHECK:   call void @_ZN2UCC1Ev(%struct.UC* %uc)
   UD ud;
+  // undefined destructor -- not allowed
   // CHECK-NOT: call
   ECI eci;
+  // empty constructor w/ initializer list -- not allowed
   // CHECK:   call void @_ZN3ECIC1Ev(%struct.ECI* %eci)
   NEC nec;
+  // non-empty constructor -- not allowed
   // CHECK:   call void @_ZN3NECC1Ev(%struct.NEC* %nec)
+  // non-empty destructor -- not allowed
   NED ned;
+  // no-constructor,  virtual method -- not allowed
   // CHECK:   call void @_ZN3NCVC1Ev(%struct.NCV* %ncv)
   NCV ncv;
   // CHECK-NOT: call