[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.

lib/CodeGen/CodeGenModule.cpp

     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 (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>() &&
+      // 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()) &&
           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)

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]])