[CUDA] Allow local `static const {__constant__, __device__}` variables.

[tra]

Summary

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

[jlebar]

wha... As you know, const doesn't mean anything, that can be const-casted away. And then you'll be able to observe that this nominally-static variable is just a normal variable.

Since this doesn't make sense and contradicts their documentation, I'm tempted to say this should only apply to the nvidia headers. Is that technically possible? And then we file a bug against nvidia and/or ask Bryce?

[yaxunl]

clang/test/CodeGenCUDA/static-device-var-no-rdc.cu

84

what happens to a const static device or constant var with non-trivial constructor? can we have a test for that?

[tra]

In D88345#2296118, @jlebar wrote:

wha... As you know, const doesn't mean anything, that can be const-casted away. And then you'll be able to observe that this nominally-static variable is just a normal variable.

Yes, I'm aware of that. It is irrelevant in this case.
The point is that __constant__ and __device__ are not allowed for the local static variables. I believe originally only __shared__ was allowed for local static vars is because __shared__ is already implicitly static, while __constant__ and __device__ are not.

On the other hand, nothing stops us giving static __constant__ of __device__ variables static storage class and that's what NVIDIA has apparently done, though without updating the docs.
The checks on the constructor still apply -- trivialness is enforced, const must have an initializer, etc.

Looks like the const-ness check should not be there, either. NVCC allows non-const statics. We may get rid of this check altogether and allow everything.

Since this doesn't make sense and contradicts their documentation, I'm tempted to say this should only apply to the nvidia headers. Is that technically possible? And then we file a bug against nvidia and/or ask Bryce?

IMO the relaxation is sensible. A static var is just a global var in an odd namespace and we do support that.
E.g. these produce the same PTX for c (modulo mangling of the names):

namespace a {
  __constant__ int c = 4;
}
__device__ void foo() {
  static __constant__ int c = 4;
}

clang/test/CodeGenCUDA/static-device-var-no-rdc.cu

84

I believe constructor trivialness check is orthogonal and will still be applied.
I'll add a test.

[jlebar]

OK, backing up, what are the semantics of static on __constant__, __device__, and __shared__?

• My understanding is that __shared__ behaves the same whether or not it's static. It's not equivalent to namespace a { __shared__ int c = 4; }, because that's illegal.
• Does __constant__ behave the same whether or not it's static? A static __constant__ is equivalent to namespace a { __constant__ int c = 4; }, and a non-static __constant__ is *also* equivalent to that?
• And __device__ does not behave the same whether or not it's static? In function scope __device__ int x; is a variable local to the thread. Whereas in global scope __device__ int x; is a global variable that lives in device memory (?). In function scope static __device__ int x; is equivalent to static int x; which is equivalent to int x; in namespace scope?

Should we mandate that you initialize static __constant__ variables in function scope? That is, if you write static __constant__ int x; in a function, then x is always uninitialized (right)? You should do static __constant__ int x = 42;?

[tra]

In D88345#2298688, @jlebar wrote:

OK, backing up, what are the semantics of static on __constant__, __device__, and __shared__?

• My understanding is that __shared__ behaves the same whether or not it's static. It's not equivalent to namespace a { __shared__ int c = 4; }, because that's illegal.

Yes. __shared__ is an odd duck. It is implicitly static, so whether we explicitly specify static makes no difference.
We're not changing anything about how it's implemented.

• Does __constant__ behave the same whether or not it's static? A static __constant__ is equivalent to namespace a { __constant__ int c = 4; }, and a non-static __constant__ is *also* equivalent to that?

No. __constant__ is not allowed on non-static local variables as it can't be allocated on stack.

• And __device__ does not behave the same whether or not it's static?

Correct.

In function scope __device__ int x; is a variable local to the thread.

Correct. __device__ in a device function is effectively a no-op and can be placed on stack as a regular local variable.

Whereas in global scope __device__ int x; is a global variable that lives in device memory (?).

Correct.

In function scope static __device__ int x; is equivalent to static int x; which is equivalent to int x; in namespace scope?

Yes, assuming you mean a __device__ function and __device__ int x; in the namespace scope.

Should we mandate that you initialize static __constant__ variables in function scope?
That is, if you write static __constant__ int x; in a function, then x is always uninitialized (right)? You should do static __constant__ int x = 42;?

No. Accoring to PTX spec: Variables in .const and .global state spaces are initialized to zero by default.
Those are the address spaces __constant__ and __device__ variables map to.
https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#state-spaces

[jlebar]

OK, now I'm starting to I understand this change..

Before, in function scope, we allow static const/non-const __shared__, and allow static const so long as it's not __device__ or __constant__.

• static -> error? (I understood us saying above that it is, but now that I read the code, isn't it saying it's an error?)
• static const -> allowed
• static __device__ -> error
• static const __device__ -> error
• static __constant__ -> error
• static const __constant__ -> error

After, in function scope, the rule is, allow static const/non-const __shared__ or anything that's static const.

• static -> error, must be const
• static const -> allowed
• static __device__ -> error, must be const
• static const __device__ -> allowed
• static __constant__ -> error, must be const
• static const __constant__ -> allowed

I guess my question when I write out this table is, why shouldn't it be like this?

• static -> allowed
• static const -> allowed
• static __device__ -> allowed
• const static __device__ -> allowed
• static __constant__ -> error, must be const
• const static __constant__ -> allowed

This makes some sense to me because we're saying, "__constant__ must be const", otherwise, anything goes.

Or here's another way of thinking about it. You're saying that static and static __device__ in function scope are the same as a __device__ variable in block scope. And a __device__ variable in block scope doesn't have to be const (right?). So why the extra restriction on function-scope static?

clang/lib/Sema/SemaDecl.cpp

13177–13178

Update comment?

[tra]

In D88345#2298879, @jlebar wrote:

OK, now I'm starting to I understand this change..

Before, in function scope, we allow static const/non-const __shared__, and allow static const so long as it's not __device__ or __constant__.

• static -> error? (I understood us saying above that it is, but now that I read the code, isn't it saying it's an error?)
• static const -> allowed
• static __device__ -> error
• static const __device__ -> error
• static __constant__ -> error
• static const __constant__ -> error

After, in function scope, the rule is, allow static const/non-const __shared__ or anything that's static const.

• static -> error, must be const
• static const -> allowed
• static __device__ -> error, must be const
• static const __device__ -> allowed
• static __constant__ -> error, must be const
• static const __constant__ -> allowed

I guess my question when I write out this table is, why shouldn't it be like this?

• static -> allowed
• static const -> allowed
• static __device__ -> allowed
• const static __device__ -> allowed
• static __constant__ -> error, must be const
• const static __constant__ -> allowed

It should. I did mention in a previous comment that > Looks like the const-ness check should not be there, either.
I need to revise the patch.

This makes some sense to me because we're saying, "__constant__ must be const", otherwise, anything goes.

Except that NVCC allows non-const __constant__, too. Generally speaking, C++ does not care about the attributes. While technically __constant__ is not changeable from the device code, not specifying const is a missed optimization/diagnostic opportunity, but not an error per se. It does not affect how the variable is emitted, but rather what user can do with it and that's beyond the scope of this patch. I don't think it warrants a hard error. A warning, perhaps, that non-const __constant__ is probably an error?

Or here's another way of thinking about it. You're saying that static and static __device__ in function scope are the same as a __device__ variable in block scope. And a __device__ variable in block scope doesn't have to be const (right?). So why the extra restriction on function-scope static?

Something like that.

[jlebar]

It should. I did mention in a previous comment that > Looks like the const-ness check should not be there, either. I need to revise the patch.

Heh, okay. Sorry I missed that, somehow this patch was confusing to me.

Except that NVCC allows non-const constant, too. Generally speaking, C++ does not care about the attributes. While technically constant is not changeable from the device code, not specifying const is a missed optimization/diagnostic opportunity, but not an error per se. It does not affect how the variable is emitted, but rather what user can do with it and that's beyond the scope of this patch. I don't think it warrants a hard error. A warning, perhaps, that non-const constant is probably an error?

Sure, that makes sense to me.

[tra]

Further relaxed application of attributes on local static variables.

[tra]

Fixed a test.

[yaxunl]

clang/test/SemaCUDA/device-var-init.cu

416–417

how does this work in device compilation? Is this equivalent to static __device__ int x = 42?

[tra]

I've verified that clang with this patch can compile Tensorflow and that it can also compile cooperative_groups.h from CUDA-11.

clang/test/SemaCUDA/device-var-init.cu

416–417

Correct.

[yaxunl]

clang/test/SemaCUDA/device-var-init.cu

416–417

so static variable without __device__/__constant__ attribute in host device function implies __device__ attribute in device compilation.

Is this also true in device function? We need Sema and CodeGen tests for these cases.

Also, can we document these changes? It is easily forgotten.

[tra]

clang/test/SemaCUDA/device-var-init.cu

416–417

I think of it as a static variable in a __device__ function. There should be no host-side shadow for it, which would normally be created for a __device__ variable.

The tests at the beginning of df_sema() in SemaCUDA/device-var-init.cu already check that static w/o attributes is allowed in __device__.
I'll update CodeGen tests to verify that we generate correct code.

[jlebar]

Hey, I'm leaving on a vacation tomorrow and didn't have a chance to get to
this review today.

Is that ok? I'm not bringing my work laptop, but I could look at it on my
personal laptop.

[jlebar]

clang/include/clang/Basic/DiagnosticSemaKinds.td

8189

__device__ is not allowed on non-static function-scope variables?

This appears to be more restrictive than we were before. I want to check, are we OK with the possibility that this will break user code? https://gcc.godbolt.org/z/Y85GKe work with clang, though not with nvcc.

I notice that we even allow __device__ int x; in __host__ __device__ functions, which is...questionable. :) https://gcc.godbolt.org/z/GjjMGx

I'm OK matching the nvcc behavior here and accepting user breakage so long as we're intentional about it. Possibly should be called out in relnotes?

clang/lib/Sema/SemaDeclAttr.cpp

4397–4398

So just to check, in our new world, __constant__ variables don't have to be const. That matches nvcc, fine.

[tra]

Added few test cases for allowed initializers on device-side static vars.

[tra]

clang/include/clang/Basic/DiagnosticSemaKinds.td

8189

It appears to have been an oversight. AFAICT, we just ignored the __device__ attribute of the local vars that sneaked past the isStaticLocal() check.

clang/lib/Sema/SemaDeclAttr.cpp

4397–4398

No. IMO, it's similar to explicitly putting a non-const variable into a .rodata section -- inadvisable, probably not very useful, and possibly implementation-defined, but not illegal.

clang/test/CodeGenCUDA/static-device-var-no-rdc.cu

84

This is checked in SemaCUDA/device-var-init.cu.

[tra]

Remove the assert which is no longer valid.

clang/include/clang/Basic/DiagnosticSemaKinds.td

 def err_cuda_device_exceptions : Error<
   "cannot use '%0' in "
   "%select{__device__|__global__|__host__|__host__ __device__}1 function">;
 def err_dynamic_var_init : Error<
     "dynamic initialization is not supported for "
     "__device__, __constant__, and __shared__ variables.">;
 def err_shared_var_init : Error<
     "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'">;
 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_host_shared : Error<
     "__shared__ local variables not allowed in "
     "%select{__device__|__global__|__host__|__host__ __device__}0 functions">;
-def err_cuda_nonglobal_constant : Error<"__constant__ variables must be global">;
+def err_cuda_nonstatic_constdev: Error<"__constant__ and __device__ are not allowed on non-static local variables">;
 def err_cuda_ovl_target : Error<
   "%select{__device__|__global__|__host__|__host__ __device__}0 function %1 "
   "cannot overload %select{__device__|__global__|__host__|__host__ __device__}2 function %3">;
 def note_cuda_ovl_candidate_target_mismatch : Note<
     "candidate template ignored: target attributes do not match">;
 
 def err_cuda_device_builtin_surftex_cls_template : Error<
     "illegal device builtin %select{surface|texture}0 reference "

clang/lib/Sema/SemaCUDA.cpp

 void Sema::checkAllowedCUDAInitializer(VarDecl *VD) {
   if (VD->isInvalidDecl() || !VD->hasInit() || !VD->hasGlobalStorage())
     return;
   const Expr *Init = VD->getInit();
   if (VD->hasAttr<CUDADeviceAttr>() || VD->hasAttr<CUDAConstantAttr>() ||
       VD->hasAttr<CUDASharedAttr>()) {
     if (LangOpts.GPUAllowDeviceInit)
       return;
-    assert(!VD->isStaticLocal() || VD->hasAttr<CUDASharedAttr>());
     bool AllowedInit = false;
     if (const CXXConstructExpr *CE = dyn_cast<CXXConstructExpr>(Init))
       AllowedInit =
           isEmptyCudaConstructor(VD->getLocation(), CE->getConstructor());
     // We'll allow constant initializers even if it's a non-empty
     // constructor according to CUDA rules. This deviates from NVCC,
     // but allows us to handle things like constexpr constructors.
     if (!AllowedInit &&

clang/lib/Sema/SemaDecl.cpp

if (VD->getTLSKind() && !hasDependentAlignment(VD) &&
       if (Context.getDeclAlign(VD) > MaxAlignChars) {
         Diag(VD->getLocation(), diag::err_tls_var_aligned_over_maximum)
           << (unsigned)Context.getDeclAlign(VD).getQuantity() << VD
           << (unsigned)MaxAlignChars.getQuantity();
       }
     }
   }
 
-  if (VD->isStaticLocal()) {
+  if (VD->isStaticLocal())
     CheckStaticLocalForDllExport(VD);
 
-    if (dyn_cast_or_null<FunctionDecl>(VD->getParentFunctionOrMethod())) {
-      // 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)
-          return;
-        if (VD->hasAttr<CUDASharedAttr>())
-          return;
-        if (VD->getType().isConstQualified() &&
-            !(VD->hasAttr<CUDADeviceAttr>() || VD->hasAttr<CUDAConstantAttr>()))
-          return;
-        if (CUDADiagIfDeviceCode(VD->getLocation(),
-                                 diag::err_device_static_local_var)
-            << CurrentCUDATarget())
-          VD->setInvalidDecl();
-      }();
-    }
-  }
-
   // Perform check for initializers of device-side global variables.
   // CUDA allows empty constructors as initializers (see E.2.3.1, CUDA
   // 7.5). We must also apply the same checks to all __shared__
   // variables whether they are local or not. CUDA also allows
   // constant initializers for __constant__ and __device__ variables.
   if (getLangOpts().CUDA)
     checkAllowedCUDAInitializer(VD);
 
   // Grab the dllimport or dllexport attribute off of the VarDecl.

clang/lib/Sema/SemaDeclAttr.cpp

static void handleOptimizeNoneAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
   if (OptimizeNoneAttr *Optnone = S.mergeOptimizeNoneAttr(D, AL))
     D->addAttr(Optnone);
 }
 
 static void handleConstantAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
   if (checkAttrMutualExclusion<CUDASharedAttr>(S, D, AL))
     return;
   const auto *VD = cast<VarDecl>(D);
-  if (!VD->hasGlobalStorage()) {
-    S.Diag(AL.getLoc(), diag::err_cuda_nonglobal_constant);
+  if (VD->hasLocalStorage()) {
+    S.Diag(AL.getLoc(), diag::err_cuda_nonstatic_constdev);
     return;
   }
   D->addAttr(::new (S.Context) CUDAConstantAttr(S.Context, AL));
 }
 
 static void handleSharedAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
   if (checkAttrMutualExclusion<CUDAConstantAttr>(S, D, AL))
     return;
static void handleGlobalAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
   // In host compilation the kernel is emitted as a stub function, which is
   // a helper function for launching the kernel. The instructions in the helper
   // function has nothing to do with the source code of the kernel. Do not emit
   // debug info for the stub function to avoid confusing the debugger.
   if (S.LangOpts.HIP && !S.LangOpts.CUDAIsDevice)
     D->addAttr(NoDebugAttr::CreateImplicit(S.Context));
 }
 
+static void handleDeviceAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
+  if (checkAttrMutualExclusion<CUDAGlobalAttr>(S, D, AL)) {
+    return;
+  }
+
+  if (const auto *VD = dyn_cast<VarDecl>(D)) {
+    if (VD->hasLocalStorage()) {
+      S.Diag(AL.getLoc(), diag::err_cuda_nonstatic_constdev);
+      return;
+    }
+  }
+  D->addAttr(::new (S.Context) CUDADeviceAttr(S.Context, AL));
+}
+
 static void handleGNUInlineAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
   const auto *Fn = cast<FunctionDecl>(D);
   if (!Fn->isInlineSpecified()) {
     S.Diag(AL.getLoc(), diag::warn_gnu_inline_attribute_requires_inline);
     return;
   }
 
   if (S.LangOpts.CPlusPlus && Fn->getStorageClass() != SC_Extern)
case ParsedAttr::AT_FormatArg:
     break;
   case ParsedAttr::AT_Callback:
     handleCallbackAttr(S, D, AL);
     break;
   case ParsedAttr::AT_CUDAGlobal:
     handleGlobalAttr(S, D, AL);
     break;
   case ParsedAttr::AT_CUDADevice:
-    handleSimpleAttributeWithExclusions<CUDADeviceAttr, CUDAGlobalAttr>(S, D,
-                                                                        AL);
+    handleDeviceAttr(S, D, AL);
     break;
   case ParsedAttr::AT_CUDAHost:
     handleSimpleAttributeWithExclusions<CUDAHostAttr, CUDAGlobalAttr>(S, D, AL);
     break;
   case ParsedAttr::AT_CUDADeviceBuiltinSurfaceType:
     handleSimpleAttributeWithExclusions<CUDADeviceBuiltinSurfaceTypeAttr,
                                         CUDADeviceBuiltinTextureTypeAttr>(S, D,
                                                                           AL);

clang/test/CodeGenCUDA/static-device-var-no-rdc.cu

 // REQUIRES: x86-registered-target
 // REQUIRES: amdgpu-registered-target
 
 // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -std=c++11 \
 // RUN:   -emit-llvm -o - -x hip %s | FileCheck \
 // RUN:   -check-prefixes=DEV %s
 
 // RUN: %clang_cc1 -triple x86_64-gnu-linux -std=c++11 \
 // RUN:   -emit-llvm -o - -x hip %s | FileCheck \
 // RUN:   -check-prefixes=HOST %s
 
 #include "Inputs/cuda.h"
 
 // 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
 // HOST-DAG: @_ZL1x = internal global i32 undef
 // HOST-DAG: @[[DEVNAMEX:[0-9]+]] = {{.*}}c"_ZL1x\00"
 
 static __device__ int x;
 

 inline __device__ void devfun(const int ** b) {
   const static int p = 2;
   b[0] = &p;
   b[1] = &x2;
 }
 
 __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;
   a[3] = x3;
   a[4] = x4;
   a[5] = x5;
   b[0] = &w;
   b[1] = &z2;
+  b[2] = &local_static_constant;
+  b[3] = &local_static_device;
   devfun(b);
 }
 
 __host__ __device__ void hdf(int *a) {
   a[0] = x3;
 }
 
 int* getDeviceSymbol(int *x);

clang/test/SemaCUDA/bad-attributes.cu

 
 __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}}
+  __constant__ int c; // expected-error {{__constant__ and __device__ are not allowed on non-static local variables}}
   __shared__ int s; // expected-error {{__shared__ local variables not allowed in __host__ functions}}
 }
 __device__ void device_fn() {
-  __constant__ int c; // expected-error {{__constant__ variables must be global}}
+  __constant__ int c; // expected-error {{__constant__ and __device__ are not allowed on non-static local variables}}
 }
 
 typedef __attribute__((device_builtin_surface_type)) unsigned long long s0_ty; // expected-warning {{'device_builtin_surface_type' attribute only applies to classes}}
 typedef __attribute__((device_builtin_texture_type)) unsigned long long t0_ty; // expected-warning {{'device_builtin_texture_type' attribute only applies to classes}}
 
 struct __attribute__((device_builtin_surface_type)) s1_ref {}; // expected-error {{illegal device builtin surface reference type 's1_ref' declared here}}
 // expected-note@-1 {{'s1_ref' needs to be instantiated from a class template with proper template arguments}}
 struct __attribute__((device_builtin_texture_type)) t1_ref {}; // expected-error {{illegal device builtin texture reference type 't1_ref' declared here}}

clang/test/SemaCUDA/device-var-init.cu

 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
 __shared__ int s_v_f = f();
 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
 __constant__ int c_v_f = f();
 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
 
 __shared__ T s_t_i = {2};
 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
+__device__ T d_t_i = {2};
+__constant__ T c_t_i = {2};
+
+__device__ ECD d_ecd_i{};
+__shared__ ECD s_ecd_i{};
+__constant__ ECD c_ecd_i{};
 
 __device__ EC d_ec_i(3);
 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
 __shared__ EC s_ec_i(3);
 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
 __constant__ EC c_ec_i(3);
 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
 

 
 __device__ T_FA_NED d_t_fa_ned;
 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
 __shared__ T_FA_NED s_t_fa_ned;
 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
 __constant__ T_FA_NED c_t_fa_ned;
 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
 
-// Verify that only __shared__ local variables may be static on device
-// side and that they are not allowed to be initialized.
+// Verify that local variables may be static on device
+// side and that they conform to the initialization constraints.
+// __shared__ can't be initialized at all and others don't support dynamic initialization.
 __device__ void df_sema() {
-  static __shared__ NCFS s_ncfs;
-  // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
-  static __shared__ UC s_uc;
-  // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
-  static __shared__ NED s_ned;
-  // 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'}}
   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'}}
   static int v;
-  // expected-error@-1 {{within a __device__ function, only __shared__ variables or const variables without device memory qualifier 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'}}
+
+
+  // __shared__ does not need to be explicitly static.
+  __shared__ int lsi;
+  // __constant__ and __device__ can not be non-static local
+  __constant__ int lci;
+  // expected-error@-1 {{__constant__ and __device__ are not allowed on non-static local variables}}
+  __device__ int ldi;
+  // expected-error@-1 {{__constant__ and __device__ are not allowed on non-static local variables}}
+
+  // Same test cases as for the globals above.
+
+  static __device__ int d_v_f = f();
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  static __shared__ int s_v_f = f();
+  // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
+  static __constant__ int c_v_f = f();
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+  static __shared__ T s_t_i = {2};
+  // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
+  static __device__ T d_t_i = {2};
+  static __constant__ T c_t_i = {2};
+
+  static __device__ ECD d_ecd_i;
+  static __shared__ ECD s_ecd_i;
+  static __constant__ ECD c_ecd_i;
+
+  static __device__ EC d_ec_i(3);
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  static __shared__ EC s_ec_i(3);
+  // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
+  static __constant__ EC c_ec_i(3);
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+  static __device__ EC d_ec_i2 = {3};
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  static __shared__ EC s_ec_i2 = {3};
+  // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
+  static __constant__ EC c_ec_i2 = {3};
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+  static __device__ ETC d_etc_i(3);
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  static __shared__ ETC s_etc_i(3);
+  // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
+  static __constant__ ETC c_etc_i(3);
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+  static __device__ ETC d_etc_i2 = {3};
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  static __shared__ ETC s_etc_i2 = {3};
+  // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
+  static __constant__ ETC c_etc_i2 = {3};
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+  static __device__ UC d_uc;
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  static __shared__ UC s_uc;
+  // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
+  static __constant__ UC c_uc;
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+  static __device__ UD d_ud;
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  static __shared__ UD s_ud;
+  // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
+  static __constant__ UD c_ud;
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+  static __device__ ECI d_eci;
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  static __shared__ ECI s_eci;
+  // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
+  static __constant__ ECI c_eci;
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+  static __device__ NEC d_nec;
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  static __shared__ NEC s_nec;
+  // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
+  static __constant__ NEC c_nec;
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+  static __device__ NED d_ned;
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  static __shared__ NED s_ned;
+  // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
+  static __constant__ NED c_ned;
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+  static __device__ NCV d_ncv;
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  static __shared__ NCV s_ncv;
+  // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
+  static __constant__ NCV c_ncv;
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+  static __device__ VD d_vd;
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  static __shared__ VD s_vd;
+  // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
+  static __constant__ VD c_vd;
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+  static __device__ NCF d_ncf;
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  static __shared__ NCF s_ncf;
+  // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
+  static __constant__ NCF c_ncf;
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+  static __shared__ NCFS s_ncfs;
+  // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
+
+  static __device__ UTC d_utc;
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  static __shared__ UTC s_utc;
+  // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
+  static __constant__ UTC c_utc;
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+  static __device__ UTC d_utc_i(3);
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  static __shared__ UTC s_utc_i(3);
+  // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
+  static __constant__ UTC c_utc_i(3);
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+  static __device__ NETC d_netc;
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  static __shared__ NETC s_netc;
+  // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
+  static __constant__ NETC c_netc;
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+  static __device__ NETC d_netc_i(3);
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  static __shared__ NETC s_netc_i(3);
+  // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
+  static __constant__ NETC c_netc_i(3);
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+  static __device__ EC_I_EC1 d_ec_i_ec1;
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  static __shared__ EC_I_EC1 s_ec_i_ec1;
+  // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
+  static __constant__ EC_I_EC1 c_ec_i_ec1;
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+  static __device__ T_V_T d_t_v_t;
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  static __shared__ T_V_T s_t_v_t;
+  // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
+  static __constant__ T_V_T c_t_v_t;
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+  static __device__ T_B_NEC d_t_b_nec;
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  static __shared__ T_B_NEC s_t_b_nec;
+  // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
+  static __constant__ T_B_NEC c_t_b_nec;
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+  static __device__ T_F_NEC d_t_f_nec;
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  static __shared__ T_F_NEC s_t_f_nec;
+  // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
+  static __constant__ T_F_NEC c_t_f_nec;
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+  static __device__ T_FA_NEC d_t_fa_nec;
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  static __shared__ T_FA_NEC s_t_fa_nec;
+  // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
+  static __constant__ T_FA_NEC c_t_fa_nec;
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+  static __device__ T_B_NED d_t_b_ned;
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  static __shared__ T_B_NED s_t_b_ned;
+  // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
+  static __constant__ T_B_NED c_t_b_ned;
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+  static __device__ T_F_NED d_t_f_ned;
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  static __shared__ T_F_NED s_t_f_ned;
+  // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
+  static __constant__ T_F_NED c_t_f_ned;
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+  static __device__ T_FA_NED d_t_fa_ned;
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  static __shared__ T_FA_NED s_t_fa_ned;
+  // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
+  static __constant__ T_FA_NED c_t_fa_ned;
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
 }
 
 __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'}}
-#endif
 }
 
 inline __host__ __device__ void hd_emitted_host_only() {
   static int x = 42; // no error on device because this is never codegen'ed there.
 }
 void call_hd_emitted_host_only() { hd_emitted_host_only(); }
 
 // Verify that we also check field initializers in instantiated structs.
 struct NontrivialInitializer {