diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/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">;
diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp
--- a/clang/lib/Sema/SemaDecl.cpp
+++ b/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)
diff --git a/clang/test/CodeGenCUDA/static-device-var-no-rdc.cu b/clang/test/CodeGenCUDA/static-device-var-no-rdc.cu
--- a/clang/test/CodeGenCUDA/static-device-var-no-rdc.cu
+++ b/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);
 }
 
diff --git a/clang/test/SemaCUDA/device-var-init.cu b/clang/test/SemaCUDA/device-var-init.cu
--- a/clang/test/SemaCUDA/device-var-init.cu
+++ b/clang/test/SemaCUDA/device-var-init.cu
@@ -207,22 +207,20 @@
   // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
 
   static __device__ int ds;
-  // expected-error@-1 {{within a __device__ function, only __shared__ variables or const variables without device memory qualifier may be marked 'static'}}
+  // expected-error@-1 {{within a __device__ function, only __shared__ variables or const variables may be marked 'static'}}
   static __constant__ int dc;
-  // expected-error@-1 {{within a __device__ function, only __shared__ variables or const variables without device memory qualifier may be marked 'static'}}
+  // expected-error@-1 {{within a __device__ function, only __shared__ variables or const variables may be marked 'static'}}
   static int v;
-  // expected-error@-1 {{within a __device__ function, only __shared__ variables or const variables without device memory qualifier may be marked 'static'}}
+  // expected-error@-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@-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@-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 or const variables without device memory qualifier may be marked 'static'}}
+  // expected-error@-2 {{within a __host__ __device__ function, only __shared__ variables or const variables may be marked 'static'}}
 #endif
 }