summaryrefslogtreecommitdiffstats
diff options
context:
space:
mode:
-rw-r--r--clang/include/clang/Basic/DiagnosticSemaKinds.td1
-rw-r--r--clang/lib/Sema/SemaDeclAttr.cpp16
-rw-r--r--clang/test/SemaCUDA/bad-attributes.cu8
3 files changed, 23 insertions, 2 deletions
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index 66a5bce86c3..eb8bf915e85 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -6723,6 +6723,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 err_cuda_nonglobal_constant : Error<"__constant__ variables must be global">;
def warn_non_pod_vararg_with_format_string : Warning<
"cannot pass %select{non-POD|non-trivial}0 object of type %1 to variadic "
diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index ff63d7b1ce6..21cace432c4 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -3696,6 +3696,19 @@ static void handleOptimizeNoneAttr(Sema &S, Decl *D,
D->addAttr(Optnone);
}
+static void handleConstantAttr(Sema &S, Decl *D, const AttributeList &Attr) {
+ if (checkAttrMutualExclusion<CUDASharedAttr>(S, D, Attr.getRange(),
+ Attr.getName()))
+ return;
+ auto *VD = cast<VarDecl>(D);
+ if (!VD->hasGlobalStorage()) {
+ S.Diag(Attr.getLoc(), diag::err_cuda_nonglobal_constant);
+ return;
+ }
+ D->addAttr(::new (S.Context) CUDAConstantAttr(
+ Attr.getRange(), S.Context, Attr.getAttributeSpellingListIndex()));
+}
+
static void handleSharedAttr(Sema &S, Decl *D, const AttributeList &Attr) {
if (checkAttrMutualExclusion<CUDAConstantAttr>(S, D, Attr.getRange(),
Attr.getName()))
@@ -5541,8 +5554,7 @@ static void ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D,
handleCommonAttr(S, D, Attr);
break;
case AttributeList::AT_CUDAConstant:
- handleSimpleAttributeWithExclusions<CUDAConstantAttr, CUDASharedAttr>(S, D,
- Attr);
+ handleConstantAttr(S, D, Attr);
break;
case AttributeList::AT_PassObjectSize:
handlePassObjectSizeAttr(S, D, Attr);
diff --git a/clang/test/SemaCUDA/bad-attributes.cu b/clang/test/SemaCUDA/bad-attributes.cu
index f3e1688462d..4fb584aec4c 100644
--- a/clang/test/SemaCUDA/bad-attributes.cu
+++ b/clang/test/SemaCUDA/bad-attributes.cu
@@ -61,3 +61,11 @@ __global__ static inline void foobar() {};
#ifdef EXPECT_INLINE_WARNING
// expected-warning@-2 {{ignored 'inline' attribute on kernel function 'foobar'}}
#endif
+
+__constant__ int global_constant;
+void host_fn() {
+ __constant__ int c; // expected-error {{__constant__ variables must be global}}
+}
+__device__ void device_fn() {
+ __constant__ int c; // expected-error {{__constant__ variables must be global}}
+}
OpenPOWER on IntegriCloud