440 lines
25 KiB
Plaintext
440 lines
25 KiB
Plaintext
|
// REQUIRES: nvptx-registered-target
|
||
|
|
|
||
|
|
// Make sure we don't allow dynamic initialization for device
|
||
|
|
// variables, but accept empty constructors allowed by CUDA.
|
||
|
|
|
||
|
|
// RUN: %clang_cc1 -verify %s -triple nvptx64-nvidia-cuda -fcuda-is-device -std=c++11 %s
|
||
|
|
|
||
|
|
#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.}}
|
||
|
|
|
||
|
|
__device__ int d_v_f = f();
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
__shared__ int s_v_f = f();
|
||
|
|
// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
|
||
|
|
__constant__ int c_v_f = f();
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
|
||
|
|
__shared__ T s_t_i = {2};
|
||
|
|
// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
|
||
|
|
__device__ T d_t_i = {2};
|
||
|
|
__constant__ T c_t_i = {2};
|
||
|
|
|
||
|
|
__device__ ECD d_ecd_i{};
|
||
|
|
__shared__ ECD s_ecd_i{};
|
||
|
|
__constant__ ECD c_ecd_i{};
|
||
|
|
|
||
|
|
__device__ EC d_ec_i(3);
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
__shared__ EC s_ec_i(3);
|
||
|
|
// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
|
||
|
|
__constant__ EC c_ec_i(3);
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
|
||
|
|
__device__ EC d_ec_i2 = {3};
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
__shared__ EC s_ec_i2 = {3};
|
||
|
|
// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
|
||
|
|
__constant__ EC c_ec_i2 = {3};
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
|
||
|
|
__device__ ETC d_etc_i(3);
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
__shared__ ETC s_etc_i(3);
|
||
|
|
// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
|
||
|
|
__constant__ ETC c_etc_i(3);
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
|
||
|
|
__device__ ETC d_etc_i2 = {3};
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
__shared__ ETC s_etc_i2 = {3};
|
||
|
|
// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
|
||
|
|
__constant__ ETC c_etc_i2 = {3};
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
|
||
|
|
__device__ UC d_uc;
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
__shared__ UC s_uc;
|
||
|
|
// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
|
||
|
|
__constant__ UC c_uc;
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
|
||
|
|
__device__ UD d_ud;
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
__shared__ UD s_ud;
|
||
|
|
// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
|
||
|
|
__constant__ UD c_ud;
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
|
||
|
|
__device__ ECI d_eci;
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
__shared__ ECI s_eci;
|
||
|
|
// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
|
||
|
|
__constant__ ECI c_eci;
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
|
||
|
|
__device__ NEC d_nec;
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
__shared__ NEC s_nec;
|
||
|
|
// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
|
||
|
|
__constant__ NEC c_nec;
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
|
||
|
|
__device__ NED d_ned;
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
__shared__ NED s_ned;
|
||
|
|
// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
|
||
|
|
__constant__ NED c_ned;
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
|
||
|
|
__device__ NCV d_ncv;
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
__shared__ NCV s_ncv;
|
||
|
|
// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
|
||
|
|
__constant__ NCV c_ncv;
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
|
||
|
|
__device__ VD d_vd;
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
__shared__ VD s_vd;
|
||
|
|
// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
|
||
|
|
__constant__ VD c_vd;
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
|
||
|
|
__device__ NCF d_ncf;
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
__shared__ NCF s_ncf;
|
||
|
|
// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
|
||
|
|
__constant__ NCF c_ncf;
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
|
||
|
|
__shared__ NCFS s_ncfs;
|
||
|
|
// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
|
||
|
|
|
||
|
|
__device__ UTC d_utc;
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
__shared__ UTC s_utc;
|
||
|
|
// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
|
||
|
|
__constant__ UTC c_utc;
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
|
||
|
|
__device__ UTC d_utc_i(3);
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
__shared__ UTC s_utc_i(3);
|
||
|
|
// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
|
||
|
|
__constant__ UTC c_utc_i(3);
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
|
||
|
|
__device__ NETC d_netc;
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
__shared__ NETC s_netc;
|
||
|
|
// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
|
||
|
|
__constant__ NETC c_netc;
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
|
||
|
|
__device__ NETC d_netc_i(3);
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
__shared__ NETC s_netc_i(3);
|
||
|
|
// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
|
||
|
|
__constant__ NETC c_netc_i(3);
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
|
||
|
|
__device__ EC_I_EC1 d_ec_i_ec1;
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
__shared__ EC_I_EC1 s_ec_i_ec1;
|
||
|
|
// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
|
||
|
|
__constant__ EC_I_EC1 c_ec_i_ec1;
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
|
||
|
|
__device__ T_V_T d_t_v_t;
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
__shared__ T_V_T s_t_v_t;
|
||
|
|
// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
|
||
|
|
__constant__ T_V_T c_t_v_t;
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
|
||
|
|
__device__ T_B_NEC d_t_b_nec;
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
__shared__ T_B_NEC s_t_b_nec;
|
||
|
|
// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
|
||
|
|
__constant__ T_B_NEC c_t_b_nec;
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
|
||
|
|
__device__ T_F_NEC d_t_f_nec;
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
__shared__ T_F_NEC s_t_f_nec;
|
||
|
|
// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
|
||
|
|
__constant__ T_F_NEC c_t_f_nec;
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
|
||
|
|
__device__ T_FA_NEC d_t_fa_nec;
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
__shared__ T_FA_NEC s_t_fa_nec;
|
||
|
|
// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
|
||
|
|
__constant__ T_FA_NEC c_t_fa_nec;
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
|
||
|
|
__device__ T_B_NED d_t_b_ned;
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
__shared__ T_B_NED s_t_b_ned;
|
||
|
|
// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
|
||
|
|
__constant__ T_B_NED c_t_b_ned;
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
|
||
|
|
__device__ T_F_NED d_t_f_ned;
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
__shared__ T_F_NED s_t_f_ned;
|
||
|
|
// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
|
||
|
|
__constant__ T_F_NED c_t_f_ned;
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
|
||
|
|
__device__ T_FA_NED d_t_fa_ned;
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
__shared__ T_FA_NED s_t_fa_ned;
|
||
|
|
// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
|
||
|
|
__constant__ T_FA_NED c_t_fa_ned;
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
|
||
|
|
// Verify that local variables may be static on device
|
||
|
|
// side and that they conform to the initialization constraints.
|
||
|
|
// __shared__ can't be initialized at all and others don't support dynamic initialization.
|
||
|
|
__device__ void df_sema() {
|
||
|
|
static __device__ int ds;
|
||
|
|
static __constant__ int dc;
|
||
|
|
static int v;
|
||
|
|
static const int cv = 1;
|
||
|
|
static const __device__ int cds = 1;
|
||
|
|
static const __constant__ int cdc = 1;
|
||
|
|
|
||
|
|
|
||
|
|
// __shared__ does not need to be explicitly static.
|
||
|
|
__shared__ int lsi;
|
||
|
|
// __constant__, __device__, and __managed__ can not be non-static local
|
||
|
|
__constant__ int lci;
|
||
|
|
// expected-error@-1 {{__constant__, __device__, and __managed__ are not allowed on non-static local variables}}
|
||
|
|
__device__ int ldi;
|
||
|
|
// expected-error@-1 {{__constant__, __device__, and __managed__ are not allowed on non-static local variables}}
|
||
|
|
|
||
|
|
// Same test cases as for the globals above.
|
||
|
|
|
||
|
|
static __device__ int d_v_f = f();
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
static __shared__ int s_v_f = f();
|
||
|
|
// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
|
||
|
|
static __constant__ int c_v_f = f();
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
|
||
|
|
static __shared__ T s_t_i = {2};
|
||
|
|
// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
|
||
|
|
static __device__ T d_t_i = {2};
|
||
|
|
static __constant__ T c_t_i = {2};
|
||
|
|
|
||
|
|
static __device__ ECD d_ecd_i;
|
||
|
|
static __shared__ ECD s_ecd_i;
|
||
|
|
static __constant__ ECD c_ecd_i;
|
||
|
|
|
||
|
|
static __device__ EC d_ec_i(3);
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
static __shared__ EC s_ec_i(3);
|
||
|
|
// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
|
||
|
|
static __constant__ EC c_ec_i(3);
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
|
||
|
|
static __device__ EC d_ec_i2 = {3};
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
static __shared__ EC s_ec_i2 = {3};
|
||
|
|
// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
|
||
|
|
static __constant__ EC c_ec_i2 = {3};
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
|
||
|
|
static __device__ ETC d_etc_i(3);
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
static __shared__ ETC s_etc_i(3);
|
||
|
|
// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
|
||
|
|
static __constant__ ETC c_etc_i(3);
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
|
||
|
|
static __device__ ETC d_etc_i2 = {3};
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
static __shared__ ETC s_etc_i2 = {3};
|
||
|
|
// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
|
||
|
|
static __constant__ ETC c_etc_i2 = {3};
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
|
||
|
|
static __device__ UC d_uc;
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
static __shared__ UC s_uc;
|
||
|
|
// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
|
||
|
|
static __constant__ UC c_uc;
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
|
||
|
|
static __device__ UD d_ud;
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
static __shared__ UD s_ud;
|
||
|
|
// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
|
||
|
|
static __constant__ UD c_ud;
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
|
||
|
|
static __device__ ECI d_eci;
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
static __shared__ ECI s_eci;
|
||
|
|
// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
|
||
|
|
static __constant__ ECI c_eci;
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
|
||
|
|
static __device__ NEC d_nec;
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
static __shared__ NEC s_nec;
|
||
|
|
// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
|
||
|
|
static __constant__ NEC c_nec;
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
|
||
|
|
static __device__ NED d_ned;
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
static __shared__ NED s_ned;
|
||
|
|
// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
|
||
|
|
static __constant__ NED c_ned;
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
|
||
|
|
static __device__ NCV d_ncv;
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
static __shared__ NCV s_ncv;
|
||
|
|
// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
|
||
|
|
static __constant__ NCV c_ncv;
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
|
||
|
|
static __device__ VD d_vd;
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
static __shared__ VD s_vd;
|
||
|
|
// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
|
||
|
|
static __constant__ VD c_vd;
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
|
||
|
|
static __device__ NCF d_ncf;
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
static __shared__ NCF s_ncf;
|
||
|
|
// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
|
||
|
|
static __constant__ NCF c_ncf;
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
|
||
|
|
static __shared__ NCFS s_ncfs;
|
||
|
|
// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
|
||
|
|
|
||
|
|
static __device__ UTC d_utc;
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
static __shared__ UTC s_utc;
|
||
|
|
// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
|
||
|
|
static __constant__ UTC c_utc;
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
|
||
|
|
static __device__ UTC d_utc_i(3);
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
static __shared__ UTC s_utc_i(3);
|
||
|
|
// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
|
||
|
|
static __constant__ UTC c_utc_i(3);
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
|
||
|
|
static __device__ NETC d_netc;
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
static __shared__ NETC s_netc;
|
||
|
|
// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
|
||
|
|
static __constant__ NETC c_netc;
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
|
||
|
|
static __device__ NETC d_netc_i(3);
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
static __shared__ NETC s_netc_i(3);
|
||
|
|
// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
|
||
|
|
static __constant__ NETC c_netc_i(3);
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
|
||
|
|
static __device__ EC_I_EC1 d_ec_i_ec1;
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
static __shared__ EC_I_EC1 s_ec_i_ec1;
|
||
|
|
// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
|
||
|
|
static __constant__ EC_I_EC1 c_ec_i_ec1;
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
|
||
|
|
static __device__ T_V_T d_t_v_t;
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
static __shared__ T_V_T s_t_v_t;
|
||
|
|
// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
|
||
|
|
static __constant__ T_V_T c_t_v_t;
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
|
||
|
|
static __device__ T_B_NEC d_t_b_nec;
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
static __shared__ T_B_NEC s_t_b_nec;
|
||
|
|
// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
|
||
|
|
static __constant__ T_B_NEC c_t_b_nec;
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
|
||
|
|
static __device__ T_F_NEC d_t_f_nec;
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
static __shared__ T_F_NEC s_t_f_nec;
|
||
|
|
// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
|
||
|
|
static __constant__ T_F_NEC c_t_f_nec;
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
|
||
|
|
static __device__ T_FA_NEC d_t_fa_nec;
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
static __shared__ T_FA_NEC s_t_fa_nec;
|
||
|
|
// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
|
||
|
|
static __constant__ T_FA_NEC c_t_fa_nec;
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
|
||
|
|
static __device__ T_B_NED d_t_b_ned;
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
static __shared__ T_B_NED s_t_b_ned;
|
||
|
|
// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
|
||
|
|
static __constant__ T_B_NED c_t_b_ned;
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
|
||
|
|
static __device__ T_F_NED d_t_f_ned;
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
static __shared__ T_F_NED s_t_f_ned;
|
||
|
|
// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
|
||
|
|
static __constant__ T_F_NED c_t_f_ned;
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
|
||
|
|
static __device__ T_FA_NED d_t_fa_ned;
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
static __shared__ T_FA_NED s_t_fa_ned;
|
||
|
|
// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
|
||
|
|
static __constant__ T_FA_NED c_t_fa_ned;
|
||
|
|
// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
|
||
|
|
}
|
||
|
|
|
||
|
|
__host__ __device__ void hd_sema() {
|
||
|
|
static int x = 42;
|
||
|
|
}
|
||
|
|
|
||
|
|
inline __host__ __device__ void hd_emitted_host_only() {
|
||
|
|
static int x = 42; // no error on device because this is never codegen'ed there.
|
||
|
|
}
|
||
|
|
void call_hd_emitted_host_only() { hd_emitted_host_only(); }
|
||
|
|
|
||
|
|
// Verify that we also check field initializers in instantiated structs.
|
||
|
|
struct NontrivialInitializer {
|
||
|
|
__host__ __device__ NontrivialInitializer() : x(43) {}
|
||
|
|
int x;
|
||
|
|
};
|
||
|
|
|
||
|
|
template <typename T>
|
||
|
|
__global__ void bar() {
|
||
|
|
__shared__ T bad;
|
||
|
|
// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
|
||
|
|
}
|
||
|
|
|
||
|
|
void instantiate() {
|
||
|
|
bar<NontrivialInitializer><<<1, 1>>>();
|
||
|
|
// expected-note@-1 {{in instantiation of function template specialization 'bar<NontrivialInitializer>' requested here}}
|
||
|
|
}
|