[clang] [CIR] Upstream support for generating global ctor regions (PR #161298)

[Henrich Lauko via cfe-commits]

================
@@ -38,3 +100,93 @@ cir::FuncOp CIRGenModule::codegenCXXStructor(GlobalDecl gd) {
   assert(!cir::MissingFeatures::opFuncAttributesForDefinition());
   return fn;
 }
+
+// Global variables requiring non-trivial initialization are handled
+// differently in CIR than in classic codegen. Classic codegen emits
+// a global init function (__cxx_global_var_init) and inserts
+// initialization for each global there. In CIR, we attach a ctor
+// region to the global variable and insert the initialization code
+// into the ctor region. This will be moved into the
+// __cxx_global_var_init function during the LoweringPrepare pass.
+void CIRGenModule::emitCXXGlobalVarDeclInit(const VarDecl *varDecl,
+                                            cir::GlobalOp addr,
+                                            bool performInit) {
+  QualType ty = varDecl->getType();
+
+  // TODO: handle address space
+  // The address space of a static local variable (DeclPtr) may be different
+  // from the address space of the "this" argument of the constructor. In that
+  // case, we need an addrspacecast before calling the constructor.
+  //
+  // struct StructWithCtor {
+  //   __device__ StructWithCtor() {...}
+  // };
+  // __device__ void foo() {
+  //   __shared__ StructWithCtor s;
+  //   ...
+  // }
+  //
+  // For example, in the above CUDA code, the static local variable s has a
+  // "shared" address space qualifier, but the constructor of StructWithCtor
+  // expects "this" in the "generic" address space.
+  assert(!cir::MissingFeatures::addressSpace());
+
+  // Create a CIRGenFunction to emit the initializer. While this isn't a true
+  // function, the handling works the same way.
+  CIRGenFunction cgf{*this, builder, true};
+  llvm::SaveAndRestore<CIRGenFunction *> savedCGF(curCGF, &cgf);
+  curCGF->curFn = addr;
+
+  CIRGenFunction::SourceLocRAIIObject fnLoc{cgf,
+                                            getLoc(varDecl->getLocation())};
+
+  assert(!cir::MissingFeatures::astVarDeclInterface());
+
+  if (!ty->isReferenceType()) {
+    assert(!cir::MissingFeatures::openMP());
+
+    bool needsDtor = varDecl->needsDestruction(getASTContext()) ==
+                     QualType::DK_cxx_destructor;
+    // PerformInit, constant store invariant / destroy handled below.
+    if (performInit) {
+      mlir::OpBuilder::InsertionGuard guard(builder);
+      auto *block = builder.createBlock(&addr.getCtorRegion());
+      CIRGenFunction::LexicalScope lexScope{*curCGF, addr.getLoc(),
+                                            builder.getInsertionBlock()};
+      lexScope.setAsGlobalInit();
+
+      builder.setInsertionPointToStart(block);
+      Address declAddr(getAddrOfGlobalVar(varDecl),
+                       getASTContext().getDeclAlign(varDecl));
+      emitDeclInit(cgf, varDecl, declAddr);
+      builder.setInsertionPointToEnd(block);
+      builder.create<cir::YieldOp>(addr->getLoc());
+    }
+
+    if (varDecl->getType().isConstantStorage(getASTContext(), true,
+                                             !needsDtor)) {
+      errorNYI(varDecl->getSourceRange(), "global with constant storage");
+    } else {
+      // If not constant storage we'll emit this regardless of NeedsDtor value.
+      mlir::OpBuilder::InsertionGuard guard(builder);
+      auto *block = builder.createBlock(&addr.getDtorRegion());
----------------
xlauko wrote:

Oh I see now that it uses different region for insertion.

Maybe I would suggest to set insertion inside `emitDeclDestroy` and `emitDeclInit` as it semantically needs to always set the correct insertion point for the global?

https://github.com/llvm/llvm-project/pull/161298