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

[Andy Kaylor 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());
----------------
andykaylor wrote:

Yeah, I think that makes sense. Near as I can tell, if we don't emit any `dtor` code we'll be creating an entry block and then erasing it, and I guess our canonicalization pass erases the region. It would be better to just not create it in the first place, and as you say, having the insertion point set where it's used will be better.

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