[PATCH] D49931: [CUDA][HIP] Allow function-scope static const variable

[Yaxun Liu via Phabricator via cfe-commits]

yaxunl updated this revision to Diff 157740.
yaxunl edited the summary of this revision.
yaxunl added a comment.

update diagnostic message.


https://reviews.llvm.org/D49931

Files:
  include/clang/Basic/DiagnosticSemaKinds.td
  lib/CodeGen/CodeGenModule.cpp
  lib/Sema/SemaDecl.cpp
  test/CodeGenCUDA/device-var-init.cu
  test/SemaCUDA/device-var-init.cu


Index: test/SemaCUDA/device-var-init.cu
===================================================================
--- test/SemaCUDA/device-var-init.cu
+++ test/SemaCUDA/device-var-init.cu
@@ -207,17 +207,17 @@
   // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
 
   static __device__ int ds;
-  // expected-error@-1 {{within a __device__ function, only __shared__ variables may be marked 'static'}}
+  // expected-error@-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@-1 {{within a __device__ function, only __shared__ variables may be marked 'static'}}
+  // expected-error@-1 {{within a __device__ function, only __shared__ variables or const variables without device memory qualifier may be marked 'static'}}
   static int v;
-  // expected-error@-1 {{within a __device__ function, only __shared__ variables may be marked 'static'}}
+  // expected-error@-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@-2 {{within a __host__ __device__ function, only __shared__ variables may be marked 'static'}}
+  // expected-error@-2 {{within a __host__ __device__ function, only __shared__ variables or const variables without device memory qualifier may be marked 'static'}}
 #endif
 }
 
Index: test/CodeGenCUDA/device-var-init.cu
===================================================================
--- test/CodeGenCUDA/device-var-init.cu
+++ test/CodeGenCUDA/device-var-init.cu
@@ -112,6 +112,9 @@
 // CHECK: @_ZZ2dfvE4s_ec = internal addrspace(3) global %struct.EC undef
 // CHECK: @_ZZ2dfvE5s_etc = internal addrspace(3) global %struct.ETC undef
 
+// CHECK: @_ZZ2dfvE11const_array = internal addrspace(4) constant [5 x i32] [i32 1, i32 2, i32 3, i32 4, i32 5]
+// CHECK: @_ZZ2dfvE9const_int = internal addrspace(4) constant i32 123
+
 // We should not emit global initializers for device-side variables.
 // CHECK-NOT: @__cxx_global_var_init
 
@@ -234,6 +237,9 @@
   static __shared__ ETC s_etc;
   // CHECK-NOT: call void @_ZN3ETCC1IJEEEDpT_(%struct.ETC* addrspacecast (%struct.ETC addrspace(3)* @_ZZ2dfvE5s_etc to %struct.ETC*))
 
+  static const int const_array[] = {1, 2, 3, 4, 5};
+  static const int const_int = 123;
+
   // anchor point separating constructors and destructors
   df(); // CHECK: call void @_Z2dfv()
 
Index: lib/Sema/SemaDecl.cpp
===================================================================
--- lib/Sema/SemaDecl.cpp
+++ lib/Sema/SemaDecl.cpp
@@ -11914,10 +11914,17 @@
         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 && !VD->hasAttr<CUDASharedAttr>() &&
+      // 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 &&
+          !(VD->hasAttr<CUDASharedAttr>() ||
+            (VD->getType().isConstQualified() &&
+             !VD->hasAttr<CUDADeviceAttr>() &&
+             !VD->hasAttr<CUDAConstantAttr>())) &&
           CUDADiagIfDeviceCode(VD->getLocation(),
                                diag::err_device_static_local_var)
               << CurrentCUDATarget())
Index: lib/CodeGen/CodeGenModule.cpp
===================================================================
--- lib/CodeGen/CodeGenModule.cpp
+++ lib/CodeGen/CodeGenModule.cpp
@@ -3176,6 +3176,10 @@
       return LangAS::cuda_constant;
     else if (D && D->hasAttr<CUDASharedAttr>())
       return LangAS::cuda_shared;
+    else if (D && D->hasAttr<CUDADeviceAttr>())
+      return LangAS::cuda_device;
+    else if (D && D->getType().isConstQualified())
+      return LangAS::cuda_constant;
     else
       return LangAS::cuda_device;
   }
Index: include/clang/Basic/DiagnosticSemaKinds.td
===================================================================
--- include/clang/Basic/DiagnosticSemaKinds.td
+++ include/clang/Basic/DiagnosticSemaKinds.td
@@ -7121,7 +7121,8 @@
     "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 may be marked 'static'">;
+    "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">;

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