--- /dev/null
+// 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];
+};
#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,
__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;
__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;
__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;
__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;
// 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;
// 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;
// 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;
// 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;
__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
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
--- /dev/null
+// CUDA struct types with interesting initialization properties.
+// Keep in sync with ../CodeGenCUDA/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];
+};
// Make sure we don't allow dynamic initialization for device
// variables, but accept empty constructors allowed by CUDA.
-// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -std=c++11 \
-// RUN: -I %S/.. -fsyntax-only -verify -o /dev/null %s
+// RUN: %clang_cc1 -verify %s -triple nvptx64-nvidia-cuda -fcuda-is-device -std=c++11 %s
-// Counterpart in CodeGen covers valid cases that pass Sema
-// checks. Here we'll only cover cases that trigger errors.
-#include "CodeGenCUDA/device-var-init.cu"
+#ifdef __clang__
+#include "Inputs/cuda.h"
+#endif
+
+// Use the types we share with CodeGen tests.
+#include "Inputs/cuda-initializers.h"
__shared__ int s_v_i = 1;
// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
projects / clang / commitdiff