Merge branch 'anom/const-var' into anom/constant-var

Author: anominos

clang/lib/CIR/CodeGen/CIRGenModule.cpp

@@ -569,15 +569,14 @@ bool CIRGenModule::shouldEmitCUDAGlobalVar(const VarDecl *global) const {
   // size and host-side address in order to provide access to
   // their device-side incarnations.
 
-  if (
-      global->getType()->isCUDADeviceBuiltinSurfaceType() ||
+  if (global->getType()->isCUDADeviceBuiltinSurfaceType() ||
       global->getType()->isCUDADeviceBuiltinTextureType()) {
     llvm_unreachable("NYI");
   }
 
   return !langOpts.CUDAIsDevice || global->hasAttr<CUDADeviceAttr>() ||
-          global->hasAttr<CUDAConstantAttr>() ||
-          global->hasAttr<CUDASharedAttr>();
+         global->hasAttr<CUDAConstantAttr>() ||
+         global->hasAttr<CUDASharedAttr>();
 }
 
 void CIRGenModule::emitGlobal(GlobalDecl gd) {
@@ -1497,7 +1496,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>())) {
       gv->setAttr(CUDAExternallyInitializedAttr::getMnemonic(),
                   CUDAExternallyInitializedAttr::get(&getMLIRContext()));
     }
@@ -1511,7 +1510,7 @@ void CIRGenModule::emitGlobalVarDefinition(const clang::VarDecl *d,
   // TODO(cir): If it is safe to mark the global 'constant', do so now.
   gv.setConstant((d->hasAttr<CUDAConstantAttr>() && langOpts.CUDAIsDevice) ||
                  (!needsGlobalCtor && !needsGlobalDtor &&
-                 isTypeConstant(d->getType(), true, true)));
+                  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/Transforms/TargetLowering/Targets/NVPTX.cpp

@@ -52,7 +52,7 @@ class NVPTXTargetLoweringInfo : public TargetLoweringInfo {
     case Kind::offload_global:
       return 1;
     case Kind::offload_constant:
-      return 2;
+      return 4;
     case Kind::offload_generic:
       return 4;
     default:

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