This revision was automatically updated to reflect the committed changes.
Closed by commit rL268962: [CUDA] Only __shared__ variables can be static local
on device side. (authored by tra).
Changed prior to commit:
http://reviews.llvm.org/D20034?vs=56610&id=56611#toc
Repository:
rL LLVM
http://reviews.llvm.org/D20034
Files:
cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
cfe/trunk/lib/Sema/SemaDecl.cpp
cfe/trunk/test/CodeGenCUDA/address-spaces.cu
cfe/trunk/test/CodeGenCUDA/device-var-init.cu
Index: cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
===================================================================
--- cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
+++ cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
@@ -6554,7 +6554,9 @@
"__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 __device__/__global__ function, "
+ "only __shared__ variables may be marked \"static\"">;
def warn_non_pod_vararg_with_format_string : Warning<
"cannot pass %select{non-POD|non-trivial}0 object of type %1 to variadic "
"%select{function|block|method|constructor}2; expected type from format "
Index: cfe/trunk/test/CodeGenCUDA/device-var-init.cu
===================================================================
--- cfe/trunk/test/CodeGenCUDA/device-var-init.cu
+++ cfe/trunk/test/CodeGenCUDA/device-var-init.cu
@@ -368,6 +368,14 @@
T_F_NEC t_f_nec;
T_FA_NEC t_fa_nec;
static __shared__ UC s_uc;
+#if ERROR_CASE
+ 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: call void @_ZN2ECC1Ev(%struct.EC* %ec)
Index: cfe/trunk/test/CodeGenCUDA/address-spaces.cu
===================================================================
--- cfe/trunk/test/CodeGenCUDA/address-spaces.cu
+++ cfe/trunk/test/CodeGenCUDA/address-spaces.cu
@@ -38,14 +38,6 @@
// CHECK: load i32, i32* addrspacecast (i32 addrspace(3)* @k to i32*)
k++;
- static int li;
- // CHECK: load i32, i32* addrspacecast (i32 addrspace(1)* @_ZZ3foovE2li to
i32*)
- li++;
-
- __constant__ int lj;
- // CHECK: load i32, i32* addrspacecast (i32 addrspace(4)* @_ZZ3foovE2lj to
i32*)
- lj++;
-
__shared__ int lk;
// CHECK: load i32, i32* addrspacecast (i32 addrspace(3)* @_ZZ3foovE2lk to
i32*)
lk++;
Index: cfe/trunk/lib/Sema/SemaDecl.cpp
===================================================================
--- cfe/trunk/lib/Sema/SemaDecl.cpp
+++ cfe/trunk/lib/Sema/SemaDecl.cpp
@@ -10391,15 +10391,24 @@
}
}
- // Static locals inherit dll attributes from their function.
if (VD->isStaticLocal()) {
if (FunctionDecl *FD =
dyn_cast_or_null<FunctionDecl>(VD->getParentFunctionOrMethod())) {
+ // Static locals inherit dll attributes from their function.
if (Attr *A = getDLLAttr(FD)) {
auto *NewAttr = cast<InheritableAttr>(A->clone(getASTContext()));
NewAttr->setInherited(true);
VD->addAttr(NewAttr);
}
+ // CUDA E.2.9.4: Within the body of a __device__ or __global__
+ // function, only __shared__ variables may be declared with
+ // static storage class.
+ if (getLangOpts().CUDA && getLangOpts().CUDAIsDevice &&
+ (FD->hasAttr<CUDADeviceAttr>() || FD->hasAttr<CUDAGlobalAttr>()) &&
+ !VD->hasAttr<CUDASharedAttr>()) {
+ Diag(VD->getLocation(), diag::err_device_static_local_var);
+ VD->setInvalidDecl();
+ }
}
}
Index: cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
===================================================================
--- cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
+++ cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
@@ -6554,7 +6554,9 @@
"__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 __device__/__global__ function, "
+ "only __shared__ variables may be marked \"static\"">;
def warn_non_pod_vararg_with_format_string : Warning<
"cannot pass %select{non-POD|non-trivial}0 object of type %1 to variadic "
"%select{function|block|method|constructor}2; expected type from format "
Index: cfe/trunk/test/CodeGenCUDA/device-var-init.cu
===================================================================
--- cfe/trunk/test/CodeGenCUDA/device-var-init.cu
+++ cfe/trunk/test/CodeGenCUDA/device-var-init.cu
@@ -368,6 +368,14 @@
T_F_NEC t_f_nec;
T_FA_NEC t_fa_nec;
static __shared__ UC s_uc;
+#if ERROR_CASE
+ 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: call void @_ZN2ECC1Ev(%struct.EC* %ec)
Index: cfe/trunk/test/CodeGenCUDA/address-spaces.cu
===================================================================
--- cfe/trunk/test/CodeGenCUDA/address-spaces.cu
+++ cfe/trunk/test/CodeGenCUDA/address-spaces.cu
@@ -38,14 +38,6 @@
// CHECK: load i32, i32* addrspacecast (i32 addrspace(3)* @k to i32*)
k++;
- static int li;
- // CHECK: load i32, i32* addrspacecast (i32 addrspace(1)* @_ZZ3foovE2li to i32*)
- li++;
-
- __constant__ int lj;
- // CHECK: load i32, i32* addrspacecast (i32 addrspace(4)* @_ZZ3foovE2lj to i32*)
- lj++;
-
__shared__ int lk;
// CHECK: load i32, i32* addrspacecast (i32 addrspace(3)* @_ZZ3foovE2lk to i32*)
lk++;
Index: cfe/trunk/lib/Sema/SemaDecl.cpp
===================================================================
--- cfe/trunk/lib/Sema/SemaDecl.cpp
+++ cfe/trunk/lib/Sema/SemaDecl.cpp
@@ -10391,15 +10391,24 @@
}
}
- // Static locals inherit dll attributes from their function.
if (VD->isStaticLocal()) {
if (FunctionDecl *FD =
dyn_cast_or_null<FunctionDecl>(VD->getParentFunctionOrMethod())) {
+ // Static locals inherit dll attributes from their function.
if (Attr *A = getDLLAttr(FD)) {
auto *NewAttr = cast<InheritableAttr>(A->clone(getASTContext()));
NewAttr->setInherited(true);
VD->addAttr(NewAttr);
}
+ // CUDA E.2.9.4: Within the body of a __device__ or __global__
+ // function, only __shared__ variables may be declared with
+ // static storage class.
+ if (getLangOpts().CUDA && getLangOpts().CUDAIsDevice &&
+ (FD->hasAttr<CUDADeviceAttr>() || FD->hasAttr<CUDAGlobalAttr>()) &&
+ !VD->hasAttr<CUDASharedAttr>()) {
+ Diag(VD->getLocation(), diag::err_device_static_local_var);
+ VD->setInvalidDecl();
+ }
}
}
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
