From 0a3ebb4d8d988e063e395621d162fa224fa4fb08 Mon Sep 17 00:00:00 2001 From: Artem Belevich Date: Mon, 2 Nov 2020 15:08:26 -0800 Subject: [PATCH] Revert "[CUDA] Allow local static variables with target attributes." This reverts commit f38a9e51178add132d2c8ae160787fb2175a48a4 Which triggered assertions. --- clang/include/clang/Basic/DiagnosticSemaKinds.td | 6 +- clang/lib/Sema/SemaDecl.cpp | 25 ++- clang/lib/Sema/SemaDeclAttr.cpp | 21 +- clang/test/CodeGenCUDA/static-device-var-no-rdc.cu | 6 - clang/test/SemaCUDA/bad-attributes.cu | 4 +- clang/test/SemaCUDA/device-var-init.cu | 224 ++------------------- 6 files changed, 52 insertions(+), 234 deletions(-) diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 83d968d..7555a5e 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -8179,6 +8179,10 @@ def err_dynamic_var_init : Error< "__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">; @@ -8186,7 +8190,7 @@ def err_cuda_extern_shared : Error<"__shared__ variable %0 cannot be 'extern'">; def err_cuda_host_shared : Error< "__shared__ local variables not allowed in " "%select{__device__|__global__|__host__|__host__ __device__}0 functions">; -def err_cuda_nonstatic_constdev: Error<"__constant__ and __device__ are not allowed on non-static local variables">; +def err_cuda_nonglobal_constant : Error<"__constant__ variables must be global">; 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/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp index 6af2c18..1dcf7a1 100644 --- a/clang/lib/Sema/SemaDecl.cpp +++ b/clang/lib/Sema/SemaDecl.cpp @@ -13172,9 +13172,32 @@ void Sema::FinalizeDeclaration(Decl *ThisDecl) { } } - 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 index 71fccc4..ce816c1 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -4394,8 +4394,8 @@ static void handleConstantAttr(Sema &S, Decl *D, const ParsedAttr &AL) { if (checkAttrMutualExclusion(S, D, AL)) return; const auto *VD = cast(D); - if (VD->hasLocalStorage()) { - S.Diag(AL.getLoc(), diag::err_cuda_nonstatic_constdev); + if (!VD->hasGlobalStorage()) { + S.Diag(AL.getLoc(), diag::err_cuda_nonglobal_constant); return; } D->addAttr(::new (S.Context) CUDAConstantAttr(S.Context, AL)); @@ -4456,20 +4456,6 @@ static void handleGlobalAttr(Sema &S, Decl *D, const ParsedAttr &AL) { 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()) { @@ -7540,7 +7526,8 @@ static void ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D, handleGlobalAttr(S, D, AL); break; case ParsedAttr::AT_CUDADevice: - handleDeviceAttr(S, D, AL); + handleSimpleAttributeWithExclusions(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 index d259d3d..9cb1c68 100644 --- a/clang/test/CodeGenCUDA/static-device-var-no-rdc.cu +++ b/clang/test/CodeGenCUDA/static-device-var-no-rdc.cu @@ -13,8 +13,6 @@ // 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 @@ -80,8 +78,6 @@ inline __device__ void devfun(const int ** b) { __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; @@ -90,8 +86,6 @@ __global__ void kernel(int *a, const int **b) { 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 index a990598..d72f744 100644 --- a/clang/test/SemaCUDA/bad-attributes.cu +++ b/clang/test/SemaCUDA/bad-attributes.cu @@ -64,11 +64,11 @@ __global__ static inline void foobar() {}; __constant__ int global_constant; void host_fn() { - __constant__ int c; // expected-error {{__constant__ and __device__ are not allowed on non-static local variables}} + __constant__ int c; // expected-error {{__constant__ variables must be global}} __shared__ int s; // expected-error {{__shared__ local variables not allowed in __host__ functions}} } __device__ void device_fn() { - __constant__ int c; // expected-error {{__constant__ and __device__ are not allowed on non-static local variables}} + __constant__ int c; // expected-error {{__constant__ variables must be global}} } 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 index 88350f5..dd5d19a 100644 --- a/clang/test/SemaCUDA/device-var-init.cu +++ b/clang/test/SemaCUDA/device-var-init.cu @@ -24,12 +24,6 @@ __constant__ int c_v_f = f(); __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.}} @@ -202,218 +196,34 @@ __shared__ T_FA_NED s_t_fa_ned; __constant__ T_FA_NED c_t_fa_ned; // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ 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. +// Verify that only __shared__ local variables may be static on device +// side and that they are not allowed to be initialized. __device__ void df_sema() { - static __device__ int ds; - static __constant__ int dc; - static int v; - static const int cv = 1; - static const __device__ int cds = 1; - static const __constant__ int cdc = 1; - - - // __shared__ does not need to be explicitly static. - __shared__ int lsi; - // __constant__ 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}; + static __shared__ NCFS s_ncfs; // 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.}} + 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'}} } __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() { -- 2.7.4