[CIR][CUDA] implement cuda constant variables (#1444)

This PR implements \_\_constant\_\_ variables.

#1438 only implements \_\_device\_\_ and \_\_shared\_\_ variables, 

~~This PR depends on #1445~~

Author: anominos

clang/lib/CIR/CodeGen/CIRGenModule.cpp

@@ -569,12 +569,12 @@ bool CIRGenModule::shouldEmitCUDAGlobalVar(const VarDecl *global) const {
   // size and host-side address in order to provide access to
   // their device-side incarnations.
 
-  if (global->hasAttr<CUDAConstantAttr>() ||
-      global->getType()->isCUDADeviceBuiltinTextureType()) {
+  if (global->getType()->isCUDADeviceBuiltinTextureType()) {
     llvm_unreachable("NYI");
   }
 
   return !langOpts.CUDAIsDevice || global->hasAttr<CUDADeviceAttr>() ||
+         global->hasAttr<CUDAConstantAttr>() ||
          global->hasAttr<CUDASharedAttr>() ||
          global->getType()->isCUDADeviceBuiltinSurfaceType();
 }
@@ -1492,7 +1492,7 @@ void CIRGenModule::emitGlobalVarDefinition(const clang::VarDecl *d,
     // __shared__ variables is not marked as externally initialized,
     // because they must not be initialized.
     if (linkage != cir::GlobalLinkageKind::InternalLinkage &&
-        (d->hasAttr<CUDADeviceAttr>() ||
+        (d->hasAttr<CUDADeviceAttr>() || d->hasAttr<CUDAConstantAttr>() ||
          d->getType()->isCUDADeviceBuiltinSurfaceType())) {
       gv->setAttr(CUDAExternallyInitializedAttr::getMnemonic(),
                   CUDAExternallyInitializedAttr::get(&getMLIRContext()));
@@ -1505,8 +1505,9 @@ void CIRGenModule::emitGlobalVarDefinition(const clang::VarDecl *d,
     emitter->finalize(gv);
 
   // TODO(cir): If it is safe to mark the global 'constant', do so now.
-  gv.setConstant(!needsGlobalCtor && !needsGlobalDtor &&
-                 isTypeConstant(d->getType(), true, true));
+  gv.setConstant((d->hasAttr<CUDAConstantAttr>() && langOpts.CUDAIsDevice) ||
+                 (!needsGlobalCtor && !needsGlobalDtor &&
+                  isTypeConstant(d->getType(), true, true)));
 
   // If it is in a read-only section, mark it 'constant'.
   if (const SectionAttr *sa = d->getAttr<SectionAttr>())

clang/lib/CIR/Dialect/IR/CIRAttrs.cpp

@@ -652,13 +652,12 @@ AddressSpaceAttr::getValueFromLangAS(clang::LangAS langAS) {
   case LangAS::cuda_device:
     return Kind::offload_global;
   case LangAS::opencl_constant:
+  case LangAS::cuda_constant:
     return Kind::offload_constant;
   case LangAS::opencl_private:
     return Kind::offload_private;
   case LangAS::opencl_generic:
     return Kind::offload_generic;
-  case LangAS::cuda_constant:
-    return Kind::offload_constant;
   case LangAS::opencl_global_device:
   case LangAS::opencl_global_host:
   case LangAS::sycl_global:

clang/test/CIR/CodeGen/CUDA/addrspace-lowering.cu

@@ -14,6 +14,6 @@ __device__ int b;
 
 // LLVM-DEVICE: @b = addrspace(1) {{.*}}
 
-// __constant__ int c;
+__constant__ int c;
 
-// XFAIL-LLVM-DEVICE: @c = addrspace(4) {{.*}}
+// LLVM-DEVICE: @c = addrspace(4) {{.*}}

clang/test/CIR/CodeGen/CUDA/global-vars.cu

@@ -17,3 +17,7 @@ __device__ int a;
 __shared__ int shared;
 // CIR-DEVICE: cir.global external addrspace(offload_local) @shared = #cir.undef
 // LLVM-DEVICE: @shared = addrspace(3) global i32 undef, align 4
+
+__constant__ int b;
+// CIR-DEVICE: cir.global constant external addrspace(offload_constant) @b = #cir.int<0> : !s32i {alignment = 4 : i64, cu.externally_initialized = #cir.cu.externally_initialized}
+// LLVM-DEVICE: @b = addrspace(4) externally_initialized constant i32 0, align 4