[CUDA][HIP] Allow function-scope static const variable

[yaxunl]

Summary

CUDA 8.0 E.3.9.4 says: 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.

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, which
can be emitted as a global variable in constant address space.

Currently clang only allows function-scope static variable with
__shared__ qualifier.

This patch also allows function-scope static const variable without
device memory qualifier and emits it as a global variable in constant
address space.

[yaxunl]

update diagnostic message.

[tra]

This patch also allows function-scope static const variable without device memory qualifier and emits it as a global variable in constant address space.

What does NVCC do with local static const variables?

[yaxunl]

In D49931#1178720, @tra wrote:

This patch also allows function-scope static const variable without device memory qualifier and emits it as a global variable in constant address space.

What does NVCC do with local static const variables?

Both const and non-const static var without device memory qualifier are allowed. It emits it as a global variable in global address space.

[tra]

Looks OK overall except for the huge if below.

lib/Sema/SemaDecl.cpp

11923–11930

This is rather convoluted. It would make it somewhat more readable if we could split CUDADiagIfDeviceCode into its own if().

Or, maybe use a lambda + early exit or, perhaps even goto to break down this huge if:

[&](){
   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();
}()

[yaxunl]

lib/Sema/SemaDecl.cpp

11923–11930

The lambda looks good. Will use it. Thanks!

[yaxunl]

Revised by Artem's comments.

include/clang/Basic/DiagnosticSemaKinds.td

def err_cuda_device_exceptions : Error<
   "%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 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">;
 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">;

lib/CodeGen/CodeGenModule.cpp

if (LangOpts.OpenCL) {
     return AddrSpace;
   }
 
   if (LangOpts.CUDA && LangOpts.CUDAIsDevice) {
     if (D && D->hasAttr<CUDAConstantAttr>())
       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;
   }
 
   return getTargetCodeGenInfo().getGlobalVarAddressSpace(*this, D);
 }
 
 LangAS CodeGenModule::getStringLiteralAddressSpace() const {

lib/Sema/SemaDecl.cpp

if (VD->isStaticLocal()) {
     if (FunctionDecl *FD =
             dyn_cast_or_null<FunctionDecl>(VD->getParentFunctionOrMethod())) {
       // Static locals inherit dll attributes from their function.
       if (Attr *A = getDLLAttr(FD)) {
         auto *NewAttr = cast<InheritableAttr>(A->clone(getASTContext()));
         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>() &&
-          CUDADiagIfDeviceCode(VD->getLocation(),
+      // 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.

test/CodeGenCUDA/device-var-init.cu

 __shared__ EC_I_EC s_ec_i_ec;
 // CHECK: @s_ec_i_ec = addrspace(3) global %struct.EC_I_EC undef,
 __constant__ EC_I_EC c_ec_i_ec;
 // CHECK: @c_ec_i_ec = addrspace(4) externally_initialized global %struct.EC_I_EC zeroinitializer,
 
 // 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
 
 // Make sure that initialization restrictions do not apply to local
 // variables.
 __device__ void df() {
   // NVPTX:  %[[ec:.*]] = alloca %struct.EC
   // NVPTX:  %[[ed:.*]] = alloca %struct.ED
__device__ void df() {
   // CHECK-NOT: call
   T_FA_NED t_fa_ned;
   // CHECK-NOT: call
   static __shared__ EC s_ec;
   // CHECK-NOT: call void @_ZN2ECC1Ev(%struct.EC* addrspacecast (%struct.EC addrspace(3)* @_ZZ2dfvE4s_ec to %struct.EC*))
   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()
 
   // Verify that we only call non-empty destructors
   // CHECK-NEXT: call void @_ZN8T_FA_NEDD1Ev(%struct.T_FA_NED* %[[t_fa_ned]])
   // CHECK-NEXT: call void @_ZN7T_F_NEDD1Ev(%struct.T_F_NED* %[[t_f_ned]])
   // CHECK-NEXT: call void @_ZN7T_B_NEDD1Ev(%struct.T_B_NED* %[[t_b_ned]])
   // CHECK-NEXT: call void @_ZN2VDD1Ev(%struct.VD* %[[vd]])

test/SemaCUDA/device-var-init.cu

__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 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'}}
+  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 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
 }
 
 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(); }