[PATCH] D25129: [CUDA] Disallow __constant__ local variables.

[Justin Lebar via cfe-commits]

This revision was automatically updated to reflect the committed changes.
Closed by commit rL282986: [CUDA] Disallow __constant__ local variables. 
(authored by jlebar).

Changed prior to commit:
  https://reviews.llvm.org/D25129?vs=73139&id=73168#toc

Repository:
  rL LLVM

https://reviews.llvm.org/D25129

Files:
  cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
  cfe/trunk/lib/Sema/SemaDeclAttr.cpp
  cfe/trunk/test/SemaCUDA/bad-attributes.cu


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 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 @@
     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);
Index: cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
===================================================================
--- cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
+++ cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
@@ -6723,6 +6723,7 @@
     "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 "
Index: cfe/trunk/test/SemaCUDA/bad-attributes.cu
===================================================================
--- cfe/trunk/test/SemaCUDA/bad-attributes.cu
+++ cfe/trunk/test/SemaCUDA/bad-attributes.cu
@@ -61,3 +61,11 @@
 #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}}
+}



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 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 @@
     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);
Index: cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
===================================================================
--- cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
+++ cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
@@ -6723,6 +6723,7 @@
     "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 "
Index: cfe/trunk/test/SemaCUDA/bad-attributes.cu
===================================================================
--- cfe/trunk/test/SemaCUDA/bad-attributes.cu
+++ cfe/trunk/test/SemaCUDA/bad-attributes.cu
@@ -61,3 +61,11 @@
 #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}}
+}

_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits