This revision was automatically updated to reflect the committed changes.
Closed by commit rL282985: [CUDA] Disallow 'extern __shared__' variables.
(authored by jlebar).
Changed prior to commit:
https://reviews.llvm.org/D25125?vs=73136&id=73167#toc
Repository:
rL LLVM
https://reviews.llvm.org/D25125
Files:
cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
cfe/trunk/lib/Sema/SemaDeclAttr.cpp
cfe/trunk/test/SemaCUDA/bad-attributes.cu
cfe/trunk/test/SemaCUDA/extern-shared.cu
Index: cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
===================================================================
--- cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
+++ cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
@@ -6722,6 +6722,7 @@
def err_cuda_vla : Error<
"cannot use variable-length arrays in "
"%select{__device__|__global__|__host__|__host__ __device__}0 functions">;
+def err_cuda_extern_shared : Error<"__shared__ variable %0 cannot be
'extern'">;
def warn_non_pod_vararg_with_format_string : Warning<
"cannot pass %select{non-POD|non-trivial}0 object of type %1 to variadic "
Index: cfe/trunk/test/SemaCUDA/bad-attributes.cu
===================================================================
--- cfe/trunk/test/SemaCUDA/bad-attributes.cu
+++ cfe/trunk/test/SemaCUDA/bad-attributes.cu
@@ -42,6 +42,8 @@
__shared__ __device__ int z9;
__shared__ __constant__ int z10; // expected-error {{attributes are not
compatible}}
// expected-note@-1 {{conflicting attribute is here}}
+__constant__ __shared__ int z10a; // expected-error {{attributes are not
compatible}}
+// expected-note@-1 {{conflicting attribute is here}}
__global__ __device__ void z11(); // expected-error {{attributes are not
compatible}}
// expected-note@-1 {{conflicting attribute is here}}
Index: cfe/trunk/test/SemaCUDA/extern-shared.cu
===================================================================
--- cfe/trunk/test/SemaCUDA/extern-shared.cu
+++ cfe/trunk/test/SemaCUDA/extern-shared.cu
@@ -0,0 +1,14 @@
+// RUN: %clang_cc1 -fsyntax-only -verify %s
+// RUN: %clang_cc1 -fsyntax-only -fcuda-is-device -verify %s
+
+#include "Inputs/cuda.h"
+
+__device__ void foo() {
+ extern __shared__ int x; // expected-error {{__shared__ variable 'x' cannot
be 'extern'}}
+}
+
+__host__ __device__ void bar() {
+ extern __shared__ int x; // expected-error {{__shared__ variable 'x' cannot
be 'extern'}}
+}
+
+extern __shared__ int global; // expected-error {{__shared__ variable 'global'
cannot be 'extern'}}
Index: cfe/trunk/lib/Sema/SemaDeclAttr.cpp
===================================================================
--- cfe/trunk/lib/Sema/SemaDeclAttr.cpp
+++ cfe/trunk/lib/Sema/SemaDeclAttr.cpp
@@ -3696,6 +3696,19 @@
D->addAttr(Optnone);
}
+static void handleSharedAttr(Sema &S, Decl *D, const AttributeList &Attr) {
+ if (checkAttrMutualExclusion<CUDAConstantAttr>(S, D, Attr.getRange(),
+ Attr.getName()))
+ return;
+ auto *VD = cast<VarDecl>(D);
+ if (VD->hasExternalStorage()) {
+ S.Diag(Attr.getLoc(), diag::err_cuda_extern_shared) << VD;
+ return;
+ }
+ D->addAttr(::new (S.Context) CUDASharedAttr(
+ Attr.getRange(), S.Context, Attr.getAttributeSpellingListIndex()));
+}
+
static void handleGlobalAttr(Sema &S, Decl *D, const AttributeList &Attr) {
if (checkAttrMutualExclusion<CUDADeviceAttr>(S, D, Attr.getRange(),
Attr.getName()) ||
@@ -5639,8 +5652,7 @@
handleSimpleAttribute<NoThrowAttr>(S, D, Attr);
break;
case AttributeList::AT_CUDAShared:
- handleSimpleAttributeWithExclusions<CUDASharedAttr, CUDAConstantAttr>(S, D,
-
Attr);
+ handleSharedAttr(S, D, Attr);
break;
case AttributeList::AT_VecReturn:
handleVecReturnAttr(S, D, Attr);
Index: cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
===================================================================
--- cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
+++ cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
@@ -6722,6 +6722,7 @@
def err_cuda_vla : Error<
"cannot use variable-length arrays in "
"%select{__device__|__global__|__host__|__host__ __device__}0 functions">;
+def err_cuda_extern_shared : Error<"__shared__ variable %0 cannot be 'extern'">;
def warn_non_pod_vararg_with_format_string : Warning<
"cannot pass %select{non-POD|non-trivial}0 object of type %1 to variadic "
Index: cfe/trunk/test/SemaCUDA/bad-attributes.cu
===================================================================
--- cfe/trunk/test/SemaCUDA/bad-attributes.cu
+++ cfe/trunk/test/SemaCUDA/bad-attributes.cu
@@ -42,6 +42,8 @@
__shared__ __device__ int z9;
__shared__ __constant__ int z10; // expected-error {{attributes are not compatible}}
// expected-note@-1 {{conflicting attribute is here}}
+__constant__ __shared__ int z10a; // expected-error {{attributes are not compatible}}
+// expected-note@-1 {{conflicting attribute is here}}
__global__ __device__ void z11(); // expected-error {{attributes are not compatible}}
// expected-note@-1 {{conflicting attribute is here}}
Index: cfe/trunk/test/SemaCUDA/extern-shared.cu
===================================================================
--- cfe/trunk/test/SemaCUDA/extern-shared.cu
+++ cfe/trunk/test/SemaCUDA/extern-shared.cu
@@ -0,0 +1,14 @@
+// RUN: %clang_cc1 -fsyntax-only -verify %s
+// RUN: %clang_cc1 -fsyntax-only -fcuda-is-device -verify %s
+
+#include "Inputs/cuda.h"
+
+__device__ void foo() {
+ extern __shared__ int x; // expected-error {{__shared__ variable 'x' cannot be 'extern'}}
+}
+
+__host__ __device__ void bar() {
+ extern __shared__ int x; // expected-error {{__shared__ variable 'x' cannot be 'extern'}}
+}
+
+extern __shared__ int global; // expected-error {{__shared__ variable 'global' cannot be 'extern'}}
Index: cfe/trunk/lib/Sema/SemaDeclAttr.cpp
===================================================================
--- cfe/trunk/lib/Sema/SemaDeclAttr.cpp
+++ cfe/trunk/lib/Sema/SemaDeclAttr.cpp
@@ -3696,6 +3696,19 @@
D->addAttr(Optnone);
}
+static void handleSharedAttr(Sema &S, Decl *D, const AttributeList &Attr) {
+ if (checkAttrMutualExclusion<CUDAConstantAttr>(S, D, Attr.getRange(),
+ Attr.getName()))
+ return;
+ auto *VD = cast<VarDecl>(D);
+ if (VD->hasExternalStorage()) {
+ S.Diag(Attr.getLoc(), diag::err_cuda_extern_shared) << VD;
+ return;
+ }
+ D->addAttr(::new (S.Context) CUDASharedAttr(
+ Attr.getRange(), S.Context, Attr.getAttributeSpellingListIndex()));
+}
+
static void handleGlobalAttr(Sema &S, Decl *D, const AttributeList &Attr) {
if (checkAttrMutualExclusion<CUDADeviceAttr>(S, D, Attr.getRange(),
Attr.getName()) ||
@@ -5639,8 +5652,7 @@
handleSimpleAttribute<NoThrowAttr>(S, D, Attr);
break;
case AttributeList::AT_CUDAShared:
- handleSimpleAttributeWithExclusions<CUDASharedAttr, CUDAConstantAttr>(S, D,
- Attr);
+ handleSharedAttr(S, D, Attr);
break;
case AttributeList::AT_VecReturn:
handleVecReturnAttr(S, D, Attr);
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
