From 85b6f63f425ba7da6f644625f2c01f207cb93279 Mon Sep 17 00:00:00 2001 From: Artem Belevich Date: Thu, 19 May 2016 20:13:39 +0000 Subject: [PATCH] [CUDA] Split device-var-init.cu tests into separate Sema and CodeGen parts. Codegen tests for device-side variable initialization are subset of test cases used to verify Sema's part of the job. Including CodeGenCUDA/device-var-init.cu from SemaCUDA makes it easier to keep both sides in sync. Differential Revision: http://reviews.llvm.org/D20139 llvm-svn: 270107 --- clang/test/CodeGenCUDA/device-var-init.cu | 232 +++------------------- clang/test/SemaCUDA/device-var-init.cu | 168 ++++++++++++++++ 2 files changed, 201 insertions(+), 199 deletions(-) create mode 100644 clang/test/SemaCUDA/device-var-init.cu diff --git a/clang/test/CodeGenCUDA/device-var-init.cu b/clang/test/CodeGenCUDA/device-var-init.cu index 8cf9278c8f59..8e29bddbd7c7 100644 --- a/clang/test/CodeGenCUDA/device-var-init.cu +++ b/clang/test/CodeGenCUDA/device-var-init.cu @@ -5,8 +5,6 @@ // RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -std=c++11 \ // RUN: -fno-threadsafe-statics -emit-llvm -o - %s | FileCheck %s -// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -std=c++11 \ -// RUN: -emit-llvm -DERROR_CASE -verify -o /dev/null %s #ifdef __clang__ #include "Inputs/cuda.h" @@ -89,21 +87,6 @@ __constant__ int c_v; __device__ int d_v_i = 1; // CHECK: @d_v_i = addrspace(1) externally_initialized global i32 1, -#ifdef ERROR_CASE -__shared__ int s_v_i = 1; -// expected-error@-1 {{initialization is not supported for __shared__ variables.}} -#endif -__constant__ int c_v_i = 1; -// CHECK: @c_v_i = addrspace(4) externally_initialized global i32 1, - -#ifdef ERROR_CASE -__device__ int d_v_f = f(); -// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ 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__, and __shared__ variables.}} -#endif __device__ T d_t; // CHECK: @d_t = addrspace(1) externally_initialized global %struct.T zeroinitializer @@ -113,11 +96,7 @@ __constant__ T c_t; // CHECK: @c_t = addrspace(4) externally_initialized global %struct.T zeroinitializer, __device__ T d_t_i = {2}; -// CHECKL @d_t_i = addrspace(1) externally_initialized global %struct.T { i32 2 }, -#ifdef ERROR_CASE -__shared__ T s_t_i = {2}; -// expected-error@-1 {{initialization is not supported for __shared__ variables.}} -#endif +// CHECK: @d_t_i = addrspace(1) externally_initialized global %struct.T { i32 2 }, __constant__ T c_t_i = {2}; // CHECK: @c_t_i = addrspace(4) externally_initialized global %struct.T { i32 2 }, @@ -128,118 +107,18 @@ __shared__ EC s_ec; __constant__ EC c_ec; // CHECK: @c_ec = addrspace(4) externally_initialized global %struct.EC zeroinitializer, -#if ERROR_CASE -__device__ EC d_ec_i(3); -// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ 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__, and __shared__ variables.}} - -__device__ EC d_ec_i2 = {3}; -// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ 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__, and __shared__ variables.}} -#endif - __device__ ETC d_etc; -// CHETCK: @d_etc = addrspace(1) externally_initialized global %struct.ETC zeroinitializer, +// CHECK: @d_etc = addrspace(1) externally_initialized global %struct.ETC zeroinitializer, __shared__ ETC s_etc; -// CHETCK: @s_etc = addrspace(3) global %struct.ETC undef, +// CHECK: @s_etc = addrspace(3) global %struct.ETC undef, __constant__ ETC c_etc; -// CHETCK: @c_etc = addrspace(4) externally_initialized global %struct.ETC zeroinitializer, - -#if ERROR_CASE -__device__ ETC d_etc_i(3); -// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ 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__, and __shared__ variables.}} - -__device__ ETC d_etc_i2 = {3}; -// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ 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__, and __shared__ variables.}} - -__device__ UC d_uc; -// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ 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__, and __shared__ variables.}} - -__device__ ECI d_eci; -// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ 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__, and __shared__ variables.}} - -__device__ NEC d_nec; -// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ 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__, and __shared__ variables.}} - -__device__ NCV d_ncv; -// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ 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__, and __shared__ variables.}} - -__device__ NCF d_ncf; -// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ 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__, and __shared__ variables.}} -#endif +// CHECK: @c_etc = addrspace(4) externally_initialized global %struct.ETC zeroinitializer, __device__ NCFS d_ncfs; // CHECK: @d_ncfs = addrspace(1) externally_initialized global %struct.NCFS { i32 3 } __constant__ NCFS c_ncfs; // CHECK: @c_ncfs = addrspace(4) externally_initialized global %struct.NCFS { i32 3 } -#if ERROR_CASE -__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__, and __shared__ 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__, and __shared__ variables.}} - -__device__ UTC d_utc_i(3); -// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ 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__, and __shared__ variables.}} - -__device__ NETC d_netc; -// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ 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__, and __shared__ variables.}} - -__device__ NETC d_netc_i(3); -// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ 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__, and __shared__ variables.}} -#endif - // Regular base class -- allowed struct T_B_T : T {}; __device__ T_B_T d_t_b_t; @@ -287,28 +166,9 @@ __constant__ EC_I_EC c_ec_i_ec; struct EC_I_EC1 : EC { __device__ EC_I_EC1() : EC(1) {} }; -#if ERROR_CASE -__device__ EC_I_EC1 d_ec_i_ec1; -// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ 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__, and __shared__ variables.}} -#endif // Virtual base class -- not allowed struct T_V_T : virtual T {}; -#if ERROR_CASE -__device__ T_V_T d_t_v_t; -// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ 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__, and __shared__ variables.}} -#endif - -// Make sure that we don't allow if we inherit or incapsulate -// something with disallowed initializer. // Inherited from or incapsulated class with non-empty constructor -- // not allowed @@ -320,29 +180,6 @@ struct T_FA_NEC { NEC nec[2]; }; -#if ERROR_CASE -__device__ T_B_NEC d_t_b_nec; -// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ 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__, and __shared__ variables.}} - -__device__ T_F_NEC d_t_f_nec; -// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ 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__, and __shared__ variables.}} - -__device__ T_FA_NEC d_t_fa_nec; -// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ 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__, and __shared__ variables.}} -#endif - // We should not emit global initializers for device-side variables. // CHECK-NOT: @__cxx_global_var_init @@ -350,60 +187,57 @@ __constant__ T_FA_NEC c_t_fa_nec; // variables. __device__ void df() { T t; + // CHECK-NOT: call EC ec; + // CHECK: call void @_ZN2ECC1Ev(%struct.EC* %ec) ETC etc; + // CHECK: call void @_ZN3ETCC1IJEEEDpT_(%struct.ETC* %etc) UC uc; + // CHECK: call void @_ZN2UCC1Ev(%struct.UC* %uc) ECI eci; + // CHECK: call void @_ZN3ECIC1Ev(%struct.ECI* %eci) NEC nec; + // CHECK: call void @_ZN3NECC1Ev(%struct.NEC* %nec) NCV ncv; + // CHECK-NOT: call NCF ncf; + // CHECK: call void @_ZN3NCFC1Ev(%struct.NCF* %ncf) NCFS ncfs; + // CHECK: call void @_ZN4NCFSC1Ev(%struct.NCFS* %ncfs) UTC utc; + // CHECK: call void @_ZN3UTCC1IJEEEDpT_(%struct.UTC* %utc) NETC netc; + // CHECK: call void @_ZN4NETCC1IJEEEDpT_(%struct.NETC* %netc) T_B_T t_b_t; + // CHECK-NOT: call T_F_T t_f_t; + // CHECK-NOT: call T_FA_T t_fa_t; + // CHECK-NOT: call EC_I_EC ec_i_ec; + // CHECK: call void @_ZN7EC_I_ECC1Ev(%struct.EC_I_EC* %ec_i_ec) EC_I_EC1 ec_i_ec1; + // CHECK: call void @_ZN8EC_I_EC1C1Ev(%struct.EC_I_EC1* %ec_i_ec1) T_V_T t_v_t; + // CHECK: call void @_ZN5T_V_TC1Ev(%struct.T_V_T* %t_v_t) T_B_NEC t_b_nec; + // CHECK: call void @_ZN7T_B_NECC1Ev(%struct.T_B_NEC* %t_b_nec) T_F_NEC t_f_nec; + // CHECK: call void @_ZN7T_F_NECC1Ev(%struct.T_F_NEC* %t_f_nec) T_FA_NEC t_fa_nec; + // CHECK: call void @_ZN8T_FA_NECC1Ev(%struct.T_FA_NEC* %t_fa_nec) static __shared__ EC s_ec; + // CHECK-NOT: call void @_ZN2ECC1Ev(%struct.EC* addrspacecast (%struct.EC addrspace(3)* @_ZZ2dfvE4s_ec to %struct.EC*)) static __shared__ ETC s_etc; -#if ERROR_CASE - static __shared__ NCFS s_ncfs; - // expected-error@-1 {{initialization is not supported for __shared__ variables.}} - static __shared__ UC s_uc; - // expected-error@-1 {{initialization is not supported for __shared__ variables.}} - static __device__ int ds; - // expected-error@-1 {{Within a __device__/__global__ function, only __shared__ variables may be marked "static"}} - static __constant__ int dc; - // expected-error@-1 {{Within a __device__/__global__ function, only __shared__ variables may be marked "static"}} - static int v; - // expected-error@-1 {{Within a __device__/__global__ function, only __shared__ variables may be marked "static"}} -#endif -} + // CHECK-NOT: call void @_ZN3ETCC1IJEEEDpT_(%struct.ETC* addrspacecast (%struct.ETC addrspace(3)* @_ZZ2dfvE5s_etc to %struct.ETC*)) -// CHECK: call void @_ZN2ECC1Ev(%struct.EC* %ec) -// CHECK: call void @_ZN3ETCC1IJEEEDpT_(%struct.ETC* %etc) -// CHECK: call void @_ZN2UCC1Ev(%struct.UC* %uc) -// CHECK: call void @_ZN3ECIC1Ev(%struct.ECI* %eci) -// CHECK: call void @_ZN3NECC1Ev(%struct.NEC* %nec) -// CHECK: call void @_ZN3NCVC1Ev(%struct.NCV* %ncv) -// CHECK: call void @_ZN3NCFC1Ev(%struct.NCF* %ncf) -// CHECK: call void @_ZN4NCFSC1Ev(%struct.NCFS* %ncfs) -// CHECK: call void @_ZN3UTCC1IJEEEDpT_(%struct.UTC* %utc) -// CHECK: call void @_ZN4NETCC1IJEEEDpT_(%struct.NETC* %netc) -// CHECK: call void @_ZN7EC_I_ECC1Ev(%struct.EC_I_EC* %ec_i_ec) -// CHECK: call void @_ZN8EC_I_EC1C1Ev(%struct.EC_I_EC1* %ec_i_ec1) -// CHECK: call void @_ZN5T_V_TC1Ev(%struct.T_V_T* %t_v_t) -// CHECK: call void @_ZN7T_B_NECC1Ev(%struct.T_B_NEC* %t_b_nec) -// CHECK: call void @_ZN7T_F_NECC1Ev(%struct.T_F_NEC* %t_f_nec) -// CHECK: call void @_ZN8T_FA_NECC1Ev(%struct.T_FA_NEC* %t_fa_nec) -// CHECK-NOT: call void @_ZN2ECC1Ev(%struct.EC* addrspacecast (%struct.EC addrspace(3)* @_ZZ2dfvE4s_ec to %struct.EC*)) -// CHECK-NOT: call void @_ZN3ETCC1IJEEEDpT_(%struct.ETC* addrspacecast (%struct.ETC addrspace(3)* @_ZZ2dfvE5s_etc to %struct.ETC*)) -// CHECK: ret void + // anchor point separating constructors and destructors + df(); // CHECK: call void @_Z2dfv() + + // CHECK-NOT: call + + // CHECK: ret void +} // We should not emit global init function. // CHECK-NOT: @_GLOBAL__sub_I diff --git a/clang/test/SemaCUDA/device-var-init.cu b/clang/test/SemaCUDA/device-var-init.cu new file mode 100644 index 000000000000..c717a3346330 --- /dev/null +++ b/clang/test/SemaCUDA/device-var-init.cu @@ -0,0 +1,168 @@ +// 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 -triple nvptx64-nvidia-cuda -fcuda-is-device -std=c++11 \ +// RUN: -I %S/.. -fsyntax-only -verify -o /dev/null %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" + +__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__, and __shared__ 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__, and __shared__ variables.}} + +__shared__ T s_t_i = {2}; +// expected-error@-1 {{initialization is not supported for __shared__ variables.}} + +__device__ EC d_ec_i(3); +// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ 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__, and __shared__ variables.}} + +__device__ EC d_ec_i2 = {3}; +// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ 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__, and __shared__ variables.}} + +__device__ ETC d_etc_i(3); +// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ 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__, and __shared__ variables.}} + +__device__ ETC d_etc_i2 = {3}; +// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ 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__, and __shared__ variables.}} + +__device__ UC d_uc; +// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ 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__, and __shared__ variables.}} + +__device__ ECI d_eci; +// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ 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__, and __shared__ variables.}} + +__device__ NEC d_nec; +// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ 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__, and __shared__ variables.}} + +__device__ NCV d_ncv; +// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ 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__, and __shared__ variables.}} + +__device__ NCF d_ncf; +// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ 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__, and __shared__ 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__, and __shared__ 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__, and __shared__ variables.}} + +__device__ UTC d_utc_i(3); +// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ 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__, and __shared__ variables.}} + +__device__ NETC d_netc; +// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ 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__, and __shared__ variables.}} + +__device__ NETC d_netc_i(3); +// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ 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__, and __shared__ variables.}} + +__device__ EC_I_EC1 d_ec_i_ec1; +// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ 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__, and __shared__ variables.}} + +__device__ T_V_T d_t_v_t; +// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ 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__, and __shared__ variables.}} + +__device__ T_B_NEC d_t_b_nec; +// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ 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__, and __shared__ variables.}} + +__device__ T_F_NEC d_t_f_nec; +// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ 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__, and __shared__ variables.}} + +__device__ T_FA_NEC d_t_fa_nec; +// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ 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__, and __shared__ variables.}} + +// Make sure that initialization restrictions do not apply to local +// variables. +__device__ void df_sema() { + static __shared__ NCFS s_ncfs; + // expected-error@-1 {{initialization is not supported for __shared__ variables.}} + static __shared__ UC s_uc; + // expected-error@-1 {{initialization is not supported for __shared__ variables.}} + static __device__ int ds; + // expected-error@-1 {{Within a __device__/__global__ function, only __shared__ variables may be marked "static"}} + static __constant__ int dc; + // expected-error@-1 {{Within a __device__/__global__ function, only __shared__ variables may be marked "static"}} + static int v; + // expected-error@-1 {{Within a __device__/__global__ function, only __shared__ variables may be marked "static"}} +}