// 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__ CEEC d_ceec;
__shared__ CEEC s_ceec;
__constant__ CEEC c_ceec;
__device__ CGTC d_cgtc;
__shared__ CGTC s_cgtc;
__constant__ CGTC c_cgtc;
__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;
for (int i = 0; i < 10; i++) {
static __device__ CEEC sd_ceec;
static __shared__ CEEC ss_ceec;
static __constant__ CEEC sc_ceec;
__shared__ CEEC s_ceec;
static __device__ CGTC sd_cgtc;
static __shared__ CGTC ss_cgtc;
static __constant__ CGTC sc_cgtc;
__shared__ CGTC s_cgtc;
}
// __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}}
for (int i = 0; i < 10; i++) {
static __device__ CEEC sd_ceec;
static __shared__ CEEC ss_ceec;
static __constant__ CEEC sc_ceec;
__shared__ CEEC s_ceec;
static __device__ CGTC sd_cgtc;
static __shared__ CGTC ss_cgtc;
static __constant__ CGTC sc_cgtc;
__shared__ CGTC s_cgtc;
}
}
// Check specialization of template function.
template <>
__global__ void bar<int>() {
__shared__ NontrivialInitializer bad;
// expected-error@-1 {{initialization is not supported for __shared__ variables}}
for (int i = 0; i < 10; i++) {
static __device__ CEEC sd_ceec;
static __shared__ CEEC ss_ceec;
static __constant__ CEEC sc_ceec;
__shared__ CEEC s_ceec;
static __device__ CGTC sd_cgtc;
static __shared__ CGTC ss_cgtc;
static __constant__ CGTC sc_cgtc;
__shared__ CGTC s_cgtc;
}
}
void instantiate() {
bar<NontrivialInitializer><<<1, 1>>>();
// expected-note@-1 {{in instantiation of function template specialization 'bar<NontrivialInitializer>' requested here}}
}
Codebase Browser
llvm