diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -8179,10 +8179,6 @@ "__device__, __constant__, and __shared__ variables.">; def err_shared_var_init : Error< "initialization is not supported for __shared__ variables.">; -def err_device_static_local_var : Error< - "within a %select{__device__|__global__|__host__|__host__ __device__}0 " - "function, only __shared__ variables or const variables without device " - "memory qualifier may be marked 'static'">; def err_cuda_vla : Error< "cannot use variable-length arrays in " "%select{__device__|__global__|__host__|__host__ __device__}0 functions">; @@ -8190,7 +8186,7 @@ def err_cuda_host_shared : Error< "__shared__ local variables not allowed in " "%select{__device__|__global__|__host__|__host__ __device__}0 functions">; -def err_cuda_nonglobal_constant : Error<"__constant__ variables must be global">; +def err_cuda_nonstatic_constdev: Error<"__constant__ and __device__ are not allowed on non-static local variables">; def err_cuda_ovl_target : Error< "%select{__device__|__global__|__host__|__host__ __device__}0 function %1 " "cannot overload %select{__device__|__global__|__host__|__host__ __device__}2 function %3">; diff --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp --- a/clang/lib/Sema/SemaCUDA.cpp +++ b/clang/lib/Sema/SemaCUDA.cpp @@ -519,7 +519,6 @@ VD->hasAttr()) { if (LangOpts.GPUAllowDeviceInit) return; - assert(!VD->isStaticLocal() || VD->hasAttr()); bool AllowedInit = false; if (const CXXConstructExpr *CE = dyn_cast(Init)) AllowedInit = diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp --- a/clang/lib/Sema/SemaDecl.cpp +++ b/clang/lib/Sema/SemaDecl.cpp @@ -13172,32 +13172,9 @@ } } - if (VD->isStaticLocal()) { + if (VD->isStaticLocal()) CheckStaticLocalForDllExport(VD); - if (dyn_cast_or_null(VD->getParentFunctionOrMethod())) { - // CUDA 8.0 E.3.9.4: Within the body of a __device__ or __global__ - // function, only __shared__ variables or variables without any device - // memory qualifiers may be declared with static storage class. - // Note: It is unclear how a function-scope non-const static variable - // without device memory qualifier is implemented, therefore only static - // const variable without device memory qualifier is allowed. - [&]() { - if (!getLangOpts().CUDA) - return; - if (VD->hasAttr()) - return; - if (VD->getType().isConstQualified() && - !(VD->hasAttr() || VD->hasAttr())) - return; - if (CUDADiagIfDeviceCode(VD->getLocation(), - diag::err_device_static_local_var) - << CurrentCUDATarget()) - VD->setInvalidDecl(); - }(); - } - } - // Perform check for initializers of device-side global variables. // CUDA allows empty constructors as initializers (see E.2.3.1, CUDA // 7.5). We must also apply the same checks to all __shared__ diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -4394,8 +4394,8 @@ if (checkAttrMutualExclusion(S, D, AL)) return; const auto *VD = cast(D); - if (!VD->hasGlobalStorage()) { - S.Diag(AL.getLoc(), diag::err_cuda_nonglobal_constant); + if (VD->hasLocalStorage()) { + S.Diag(AL.getLoc(), diag::err_cuda_nonstatic_constdev); return; } D->addAttr(::new (S.Context) CUDAConstantAttr(S.Context, AL)); @@ -4456,6 +4456,20 @@ D->addAttr(NoDebugAttr::CreateImplicit(S.Context)); } +static void handleDeviceAttr(Sema &S, Decl *D, const ParsedAttr &AL) { + if (checkAttrMutualExclusion(S, D, AL)) { + return; + } + + if (const auto *VD = dyn_cast(D)) { + if (VD->hasLocalStorage()) { + S.Diag(AL.getLoc(), diag::err_cuda_nonstatic_constdev); + return; + } + } + D->addAttr(::new (S.Context) CUDADeviceAttr(S.Context, AL)); +} + static void handleGNUInlineAttr(Sema &S, Decl *D, const ParsedAttr &AL) { const auto *Fn = cast(D); if (!Fn->isInlineSpecified()) { @@ -7526,8 +7540,7 @@ handleGlobalAttr(S, D, AL); break; case ParsedAttr::AT_CUDADevice: - handleSimpleAttributeWithExclusions(S, D, - AL); + handleDeviceAttr(S, D, AL); break; case ParsedAttr::AT_CUDAHost: handleSimpleAttributeWithExclusions(S, D, AL); diff --git a/clang/test/CodeGenCUDA/static-device-var-no-rdc.cu b/clang/test/CodeGenCUDA/static-device-var-no-rdc.cu --- a/clang/test/CodeGenCUDA/static-device-var-no-rdc.cu +++ b/clang/test/CodeGenCUDA/static-device-var-no-rdc.cu @@ -13,6 +13,8 @@ // Test function scope static device variable, which should not be externalized. // DEV-DAG: @_ZZ6kernelPiPPKiE1w = internal addrspace(4) constant i32 1 +// DEV-DAG: @_ZZ6kernelPiPPKiE21local_static_constant = internal addrspace(4) constant i32 42 +// DEV-DAG: @_ZZ6kernelPiPPKiE19local_static_device = internal addrspace(1) constant i32 43 // Check a static device variable referenced by host function is externalized. // DEV-DAG: @_ZL1x = addrspace(1) externally_initialized global i32 0 @@ -78,6 +80,8 @@ __global__ void kernel(int *a, const int **b) { const static int w = 1; + const static __constant__ int local_static_constant = 42; + const static __device__ int local_static_device = 43; a[0] = x; a[1] = y; a[2] = x2; @@ -86,6 +90,8 @@ a[5] = x5; b[0] = &w; b[1] = &z2; + b[2] = &local_static_constant; + b[3] = &local_static_device; devfun(b); } diff --git a/clang/test/SemaCUDA/bad-attributes.cu b/clang/test/SemaCUDA/bad-attributes.cu --- a/clang/test/SemaCUDA/bad-attributes.cu +++ b/clang/test/SemaCUDA/bad-attributes.cu @@ -64,11 +64,11 @@ __constant__ int global_constant; void host_fn() { - __constant__ int c; // expected-error {{__constant__ variables must be global}} + __constant__ int c; // expected-error {{__constant__ and __device__ are not allowed on non-static local variables}} __shared__ int s; // expected-error {{__shared__ local variables not allowed in __host__ functions}} } __device__ void device_fn() { - __constant__ int c; // expected-error {{__constant__ variables must be global}} + __constant__ int c; // expected-error {{__constant__ and __device__ are not allowed on non-static local variables}} } typedef __attribute__((device_builtin_surface_type)) unsigned long long s0_ty; // expected-warning {{'device_builtin_surface_type' attribute only applies to classes}} diff --git a/clang/test/SemaCUDA/device-var-init.cu b/clang/test/SemaCUDA/device-var-init.cu --- a/clang/test/SemaCUDA/device-var-init.cu +++ b/clang/test/SemaCUDA/device-var-init.cu @@ -24,6 +24,12 @@ __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__, and __shared__ variables.}} @@ -196,34 +202,218 @@ __constant__ T_FA_NED c_t_fa_ned; // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}} -// Verify that only __shared__ local variables may be static on device -// side and that they are not allowed to be initialized. +// 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 __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 __shared__ NED s_ned; - // expected-error@-1 {{initialization is not supported for __shared__ variables.}} - static __device__ int ds; - // expected-error@-1 {{within a __device__ function, only __shared__ variables or const variables without device memory qualifier may be marked 'static'}} static __constant__ int dc; - // expected-error@-1 {{within a __device__ function, only __shared__ variables or const variables without device memory qualifier may be marked 'static'}} static int v; - // expected-error@-1 {{within a __device__ function, only __shared__ variables or const variables without device memory qualifier may be marked 'static'}} static const int cv = 1; static const __device__ int cds = 1; - // expected-error@-1 {{within a __device__ function, only __shared__ variables or const variables without device memory qualifier may be marked 'static'}} static const __constant__ int cdc = 1; - // expected-error@-1 {{within a __device__ function, only __shared__ variables or const variables without device memory qualifier may be marked 'static'}} + + + // __shared__ does not need to be explicitly static. + __shared__ int lsi; + // __constant__ and __device__ can not be non-static local + __constant__ int lci; + // expected-error@-1 {{__constant__ and __device__ are not allowed on non-static local variables}} + __device__ int ldi; + // expected-error@-1 {{__constant__ and __device__ 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__, and __shared__ 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__, and __shared__ 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__, and __shared__ 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__, and __shared__ variables.}} + + static __device__ EC d_ec_i2 = {3}; + // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ 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__, and __shared__ variables.}} + + static __device__ ETC d_etc_i(3); + // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ 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__, and __shared__ variables.}} + + static __device__ ETC d_etc_i2 = {3}; + // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ 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__, and __shared__ variables.}} + + static __device__ UC d_uc; + // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ 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__, and __shared__ variables.}} + + static __device__ UD d_ud; + // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ 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__, and __shared__ variables.}} + + static __device__ ECI d_eci; + // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ 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__, and __shared__ variables.}} + + static __device__ NEC d_nec; + // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ 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__, and __shared__ variables.}} + + static __device__ NED d_ned; + // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ 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__, and __shared__ variables.}} + + static __device__ NCV d_ncv; + // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ 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__, and __shared__ variables.}} + + static __device__ VD d_vd; + // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ 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__, and __shared__ variables.}} + + static __device__ NCF d_ncf; + // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ 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__, and __shared__ 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__, and __shared__ 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__, and __shared__ variables.}} + + static __device__ UTC d_utc_i(3); + // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ 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__, and __shared__ variables.}} + + static __device__ NETC d_netc; + // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ 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__, and __shared__ variables.}} + + static __device__ NETC d_netc_i(3); + // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ 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__, and __shared__ variables.}} + + static __device__ EC_I_EC1 d_ec_i_ec1; + // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ 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__, and __shared__ variables.}} + + static __device__ T_V_T d_t_v_t; + // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ 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__, and __shared__ variables.}} + + static __device__ T_B_NEC d_t_b_nec; + // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ 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__, and __shared__ variables.}} + + static __device__ T_F_NEC d_t_f_nec; + // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ 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__, and __shared__ variables.}} + + static __device__ T_FA_NEC d_t_fa_nec; + // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ 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__, and __shared__ variables.}} + + static __device__ T_B_NED d_t_b_ned; + // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ 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__, and __shared__ variables.}} + + static __device__ T_F_NED d_t_f_ned; + // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ 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__, and __shared__ variables.}} + + static __device__ T_FA_NED d_t_fa_ned; + // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ 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__, and __shared__ variables.}} } __host__ __device__ void hd_sema() { static int x = 42; -#ifdef __CUDA_ARCH__ - // expected-error@-2 {{within a __host__ __device__ function, only __shared__ variables or const variables without device memory qualifier may be marked 'static'}} -#endif } inline __host__ __device__ void hd_emitted_host_only() {