[PATCH] D88345: [CUDA] Allow `static const {__constant__, __device__}` variables.

[Artem Belevich via Phabricator via cfe-commits]

tra created this revision.
tra added reviewers: jlebar, yaxunl.
Herald added subscribers: sanjoy.google, bixia.
Herald added a project: clang.
tra requested review of this revision.

While CUDA documentation claims that such variables are not allowed[1], NVCC has
been accepting them since CUDA-10.0 [2] and some headers in CUDA-11 rely on this
working.

1. https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#static-variables-function
2. https://godbolt.org/z/zsodzc


Repository:
  rG LLVM Github Monorepo

https://reviews.llvm.org/D88345

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


Index: clang/test/SemaCUDA/device-var-init.cu
===================================================================
--- clang/test/SemaCUDA/device-var-init.cu
+++ clang/test/SemaCUDA/device-var-init.cu
@@ -207,22 +207,20 @@
   // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
 
   static __device__ int ds;
-  // expected-error at -1 {{within a __device__ function, only __shared__ variables or const variables without device memory qualifier may be marked 'static'}}
+  // expected-error at -1 {{within a __device__ function, only __shared__ variables or const variables may be marked 'static'}}
   static __constant__ int dc;
-  // expected-error at -1 {{within a __device__ function, only __shared__ variables or const variables without device memory qualifier may be marked 'static'}}
+  // expected-error at -1 {{within a __device__ function, only __shared__ variables or const variables may be marked 'static'}}
   static int v;
-  // expected-error at -1 {{within a __device__ function, only __shared__ variables or const variables without device memory qualifier may be marked 'static'}}
+  // expected-error at -1 {{within a __device__ function, only __shared__ variables or const variables may be marked 'static'}}
   static const int cv = 1;
   static const __device__ int cds = 1;
-  // expected-error at -1 {{within a __device__ function, only __shared__ variables or const variables without device memory qualifier may be marked 'static'}}
   static const __constant__ int cdc = 1;
-  // expected-error at -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 at -2 {{within a __host__ __device__ function, only __shared__ variables or const variables without device memory qualifier may be marked 'static'}}
+  // expected-error at -2 {{within a __host__ __device__ function, only __shared__ variables or const variables may be marked 'static'}}
 #endif
 }
 
Index: clang/test/CodeGenCUDA/static-device-var-no-rdc.cu
===================================================================
--- clang/test/CodeGenCUDA/static-device-var-no-rdc.cu
+++ clang/test/CodeGenCUDA/static-device-var-no-rdc.cu
@@ -13,6 +13,8 @@
 
 // Test function scope static device variable, which should not be externalized.
 // DEV-DAG: @_ZZ6kernelPiPPKiE1w = internal addrspace(4) constant i32 1
+// DEV-DAG: @_ZZ6kernelPiPPKiE21local_static_constant = internal addrspace(4) constant i32 42
+// DEV-DAG: @_ZZ6kernelPiPPKiE19local_static_device = internal addrspace(1) constant i32 43
 
 // Check a static device variable referenced by host function is externalized.
 // DEV-DAG: @_ZL1x = addrspace(1) externally_initialized global i32 0
@@ -78,6 +80,8 @@
 
 __global__ void kernel(int *a, const int **b) {
   const static int w = 1;
+  const static __constant__ int local_static_constant = 42;
+  const static __device__ int local_static_device = 43;
   a[0] = x;
   a[1] = y;
   a[2] = x2;
@@ -86,6 +90,8 @@
   a[5] = x5;
   b[0] = &w;
   b[1] = &z2;
+  b[2] = &local_static_constant;
+  b[3] = &local_static_device;
   devfun(b);
 }
 
Index: clang/lib/Sema/SemaDecl.cpp
===================================================================
--- clang/lib/Sema/SemaDecl.cpp
+++ clang/lib/Sema/SemaDecl.cpp
@@ -13162,10 +13162,7 @@
       [&]() {
         if (!getLangOpts().CUDA)
           return;
-        if (VD->hasAttr<CUDASharedAttr>())
-          return;
-        if (VD->getType().isConstQualified() &&
-            !(VD->hasAttr<CUDADeviceAttr>() || VD->hasAttr<CUDAConstantAttr>()))
+        if (VD->hasAttr<CUDASharedAttr>() || VD->getType().isConstQualified())
           return;
         if (CUDADiagIfDeviceCode(VD->getLocation(),
                                  diag::err_device_static_local_var)
Index: clang/include/clang/Basic/DiagnosticSemaKinds.td
===================================================================
--- clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -8155,8 +8155,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 or const variables without device "
-    "memory qualifier may be marked 'static'">;
+    "function, only __shared__ variables or const variables "
+    "may be marked 'static'">;
 def err_cuda_vla : Error<
     "cannot use variable-length arrays in "
     "%select{__device__|__global__|__host__|__host__ __device__}0 functions">;


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D88345.294451.patch
Type: text/x-patch
Size: 4707 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20200925/9f06b8c9/attachment.bin>