forked from OSchip/llvm-project
[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
This commit is contained in:
Artem Belevich
parent
a972d6121e
commit
85b6f63f42
|
|
@ -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
|
||||
|
|
|
|||
|
|
@ -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"}}
|
||||
}
|
||||
Loading…
Reference in New Issue