[clang] 0a3ebb4 - Revert "[CUDA] Allow local static variables with target attributes."
Artem Belevich via cfe-commits
cfe-commits at lists.llvm.org
Mon Nov 2 15:09:25 PST 2020
Author: Artem Belevich
Date: 2020-11-02T15:09:07-08:00
New Revision: 0a3ebb4d8d988e063e395621d162fa224fa4fb08
URL: https://github.com/llvm/llvm-project/commit/0a3ebb4d8d988e063e395621d162fa224fa4fb08
DIFF: https://github.com/llvm/llvm-project/commit/0a3ebb4d8d988e063e395621d162fa224fa4fb08.diff
LOG: Revert "[CUDA] Allow local static variables with target attributes."
This reverts commit f38a9e51178add132d2c8ae160787fb2175a48a4
Which triggered assertions.
Added:
Modified:
clang/include/clang/Basic/DiagnosticSemaKinds.td
clang/lib/Sema/SemaDecl.cpp
clang/lib/Sema/SemaDeclAttr.cpp
clang/test/CodeGenCUDA/static-device-var-no-rdc.cu
clang/test/SemaCUDA/bad-attributes.cu
clang/test/SemaCUDA/device-var-init.cu
Removed:
################################################################################
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index 83d968d3ae95..7555a5ebbd34 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 6af2c188ba50..1dcf7a16a897 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<FunctionDecl>(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<CUDASharedAttr>())
+ return;
+ if (VD->getType().isConstQualified() &&
+ !(VD->hasAttr<CUDADeviceAttr>() || VD->hasAttr<CUDAConstantAttr>()))
+ 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 71fccc4c68e7..ce816c1ab4b7 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<CUDASharedAttr>(S, D, AL))
return;
const auto *VD = cast<VarDecl>(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<CUDAGlobalAttr>(S, D, AL)) {
- return;
- }
-
- if (const auto *VD = dyn_cast<VarDecl>(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<FunctionDecl>(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<CUDADeviceAttr, CUDAGlobalAttr>(S, D,
+ AL);
break;
case ParsedAttr::AT_CUDAHost:
handleSimpleAttributeWithExclusions<CUDAHostAttr, CUDAGlobalAttr>(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 d259d3dd78cf..9cb1c6804a48 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 a990598e7b7e..d72f74471c23 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 88350f56651c..dd5d19a6a2e8 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 at -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 at -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 at -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 at -1 {{__constant__ and __device__ are not allowed on non-static local variables}}
- __device__ int ldi;
- // expected-error at -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 at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
- static __shared__ int s_v_f = f();
- // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
- static __constant__ int c_v_f = f();
- // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
-
- static __shared__ T s_t_i = {2};
- // expected-error at -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 at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
- static __shared__ EC s_ec_i(3);
- // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
- static __constant__ EC c_ec_i(3);
- // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
-
- static __device__ EC d_ec_i2 = {3};
- // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
- static __shared__ EC s_ec_i2 = {3};
- // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
- static __constant__ EC c_ec_i2 = {3};
- // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
-
- static __device__ ETC d_etc_i(3);
- // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
- static __shared__ ETC s_etc_i(3);
- // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
- static __constant__ ETC c_etc_i(3);
- // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
-
- static __device__ ETC d_etc_i2 = {3};
- // expected-error at -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 at -1 {{initialization is not supported for __shared__ variables.}}
- static __constant__ ETC c_etc_i2 = {3};
- // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
-
- static __device__ UC d_uc;
- // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
static __shared__ UC s_uc;
// expected-error at -1 {{initialization is not supported for __shared__ variables.}}
- static __constant__ UC c_uc;
- // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
-
- static __device__ UD d_ud;
- // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
- static __shared__ UD s_ud;
- // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
- static __constant__ UD c_ud;
- // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
-
- static __device__ ECI d_eci;
- // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
- static __shared__ ECI s_eci;
- // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
- static __constant__ ECI c_eci;
- // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
-
- static __device__ NEC d_nec;
- // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
- static __shared__ NEC s_nec;
- // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
- static __constant__ NEC c_nec;
- // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
-
- static __device__ NED d_ned;
- // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
static __shared__ NED s_ned;
// expected-error at -1 {{initialization is not supported for __shared__ variables.}}
- static __constant__ NED c_ned;
- // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
-
- static __device__ NCV d_ncv;
- // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
- static __shared__ NCV s_ncv;
- // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
- static __constant__ NCV c_ncv;
- // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
-
- static __device__ VD d_vd;
- // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
- static __shared__ VD s_vd;
- // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
- static __constant__ VD c_vd;
- // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
- static __device__ NCF d_ncf;
- // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
- static __shared__ NCF s_ncf;
- // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
- static __constant__ NCF c_ncf;
- // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
-
- static __shared__ NCFS s_ncfs;
- // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
-
- static __device__ UTC d_utc;
- // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
- static __shared__ UTC s_utc;
- // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
- static __constant__ UTC c_utc;
- // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
-
- static __device__ UTC d_utc_i(3);
- // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
- static __shared__ UTC s_utc_i(3);
- // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
- static __constant__ UTC c_utc_i(3);
- // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
-
- static __device__ NETC d_netc;
- // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
- static __shared__ NETC s_netc;
- // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
- static __constant__ NETC c_netc;
- // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
-
- static __device__ NETC d_netc_i(3);
- // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
- static __shared__ NETC s_netc_i(3);
- // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
- static __constant__ NETC c_netc_i(3);
- // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
-
- static __device__ EC_I_EC1 d_ec_i_ec1;
- // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
- static __shared__ EC_I_EC1 s_ec_i_ec1;
- // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
- static __constant__ EC_I_EC1 c_ec_i_ec1;
- // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
-
- static __device__ T_V_T d_t_v_t;
- // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
- static __shared__ T_V_T s_t_v_t;
- // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
- static __constant__ T_V_T c_t_v_t;
- // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
-
- static __device__ T_B_NEC d_t_b_nec;
- // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
- static __shared__ T_B_NEC s_t_b_nec;
- // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
- static __constant__ T_B_NEC c_t_b_nec;
- // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
-
- static __device__ T_F_NEC d_t_f_nec;
- // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
- static __shared__ T_F_NEC s_t_f_nec;
- // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
- static __constant__ T_F_NEC c_t_f_nec;
- // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
-
- static __device__ T_FA_NEC d_t_fa_nec;
- // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
- static __shared__ T_FA_NEC s_t_fa_nec;
- // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
- static __constant__ T_FA_NEC c_t_fa_nec;
- // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
-
- static __device__ T_B_NED d_t_b_ned;
- // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
- static __shared__ T_B_NED s_t_b_ned;
- // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
- static __constant__ T_B_NED c_t_b_ned;
- // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
-
- static __device__ T_F_NED d_t_f_ned;
- // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
- static __shared__ T_F_NED s_t_f_ned;
- // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
- static __constant__ T_F_NED c_t_f_ned;
- // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
-
- static __device__ T_FA_NED d_t_fa_ned;
- // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
- static __shared__ T_FA_NED s_t_fa_ned;
- // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
- static __constant__ T_FA_NED c_t_fa_ned;
- // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+ static __device__ int ds;
+ // expected-error at -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 at -1 {{within a __device__ function, only __shared__ variables or const variables without device memory qualifier may be marked 'static'}}
+ static int v;
+ // expected-error at -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 at -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 at -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 at -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() {
More information about the cfe-commits
mailing list