[clang] [CIR][OpenACC] Implement pointer/array recipe destructors (PR #160189)
via cfe-commits
cfe-commits at lists.llvm.org
Mon Sep 22 13:08:39 PDT 2025
llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-clang
Author: Erich Keane (erichkeane)
<details>
<summary>Changes</summary>
After previous implementation, I discovered that we were both doing arrays incorrectly for recipes, plus didn't get the pointer allocations done correctly. This patch is the first of a few in a series that attempts to make sure we get all pointers/arrays correct.
This patch is limited to just 'private' and destructors, which simplifies the review significantly. Destructors are simply looped through and called at each level.
The 'recipe-decl' is the 'least bounded' (that is, the type of the
expression, in the type of `int[5] i; #pragma acc parallel
private(i[1]), the type of the `recipe-decl` is `int`. This allows
us to do init/destruction at the element level.
This patch also adds infrastructure for the rest of the series of private (for the init section), as well as extensive testing for 'private', with a lot of 'TODO' locations.
Future patches will fill these in, but at the moment, there is an NYI warning for bounds, so a number of tests are updated to handle that.
---
Patch is 271.62 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/160189.diff
41 Files Affected:
- (modified) clang/include/clang/AST/OpenACCClause.h (+1-3)
- (modified) clang/lib/CIR/CodeGen/CIRGenFunction.h (+7)
- (modified) clang/lib/CIR/CodeGen/CIRGenOpenACC.cpp (+32-4)
- (modified) clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp (+21-22)
- (modified) clang/lib/CIR/CodeGen/CIRGenOpenACCRecipe.h (+255-38)
- (modified) clang/lib/Sema/SemaOpenACC.cpp (+15-4)
- (modified) clang/test/CIR/CodeGenOpenACC/combined-firstprivate-clause.cpp (+1-1)
- (modified) clang/test/CIR/CodeGenOpenACC/combined-private-clause.cpp (+68-76)
- (modified) clang/test/CIR/CodeGenOpenACC/combined-reduction-clause-default-ops.cpp (+1-1)
- (modified) clang/test/CIR/CodeGenOpenACC/combined-reduction-clause-float.cpp (+1-1)
- (modified) clang/test/CIR/CodeGenOpenACC/combined-reduction-clause-inline-ops.cpp (+1-1)
- (modified) clang/test/CIR/CodeGenOpenACC/combined-reduction-clause-int.cpp (+1-1)
- (modified) clang/test/CIR/CodeGenOpenACC/combined-reduction-clause-outline-ops.cpp (+1-1)
- (modified) clang/test/CIR/CodeGenOpenACC/compute-firstprivate-clause.c (+1-1)
- (modified) clang/test/CIR/CodeGenOpenACC/compute-firstprivate-clause.cpp (+1-1)
- (modified) clang/test/CIR/CodeGenOpenACC/compute-private-clause.c (+22-22)
- (modified) clang/test/CIR/CodeGenOpenACC/compute-private-clause.cpp (+68-76)
- (modified) clang/test/CIR/CodeGenOpenACC/compute-reduction-clause-default-ops.c (+1-1)
- (modified) clang/test/CIR/CodeGenOpenACC/compute-reduction-clause-default-ops.cpp (+1-1)
- (modified) clang/test/CIR/CodeGenOpenACC/compute-reduction-clause-float.c (+1-1)
- (modified) clang/test/CIR/CodeGenOpenACC/compute-reduction-clause-float.cpp (+1-1)
- (modified) clang/test/CIR/CodeGenOpenACC/compute-reduction-clause-inline-ops.cpp (+1-1)
- (modified) clang/test/CIR/CodeGenOpenACC/compute-reduction-clause-int.c (+1-1)
- (modified) clang/test/CIR/CodeGenOpenACC/compute-reduction-clause-int.cpp (+1-1)
- (modified) clang/test/CIR/CodeGenOpenACC/compute-reduction-clause-outline-ops.cpp (+1-1)
- (modified) clang/test/CIR/CodeGenOpenACC/compute-reduction-clause-unsigned-int.c (+1-1)
- (modified) clang/test/CIR/CodeGenOpenACC/loop-private-clause.cpp (+68-76)
- (modified) clang/test/CIR/CodeGenOpenACC/loop-reduction-clause-default-ops.cpp (+1-1)
- (modified) clang/test/CIR/CodeGenOpenACC/loop-reduction-clause-float.cpp (+1-1)
- (modified) clang/test/CIR/CodeGenOpenACC/loop-reduction-clause-inline-ops.cpp (+1-1)
- (modified) clang/test/CIR/CodeGenOpenACC/loop-reduction-clause-int.cpp (+1-1)
- (modified) clang/test/CIR/CodeGenOpenACC/loop-reduction-clause-outline-ops.cpp (+1-1)
- (added) clang/test/CIR/CodeGenOpenACC/private-clause-array-recipes-CtorDtor.cpp (+448)
- (added) clang/test/CIR/CodeGenOpenACC/private-clause-array-recipes-NoOps.cpp (+142)
- (added) clang/test/CIR/CodeGenOpenACC/private-clause-array-recipes-int.cpp (+87)
- (added) clang/test/CIR/CodeGenOpenACC/private-clause-pointer-array-recipes-CtorDtor.cpp (+845)
- (added) clang/test/CIR/CodeGenOpenACC/private-clause-pointer-array-recipes-NoOps.cpp (+253)
- (added) clang/test/CIR/CodeGenOpenACC/private-clause-pointer-array-recipes-int.cpp (+251)
- (added) clang/test/CIR/CodeGenOpenACC/private-clause-pointer-recipes-CtorDtor.cpp (+291)
- (added) clang/test/CIR/CodeGenOpenACC/private-clause-pointer-recipes-NoOps.cpp (+117)
- (added) clang/test/CIR/CodeGenOpenACC/private-clause-pointer-recipes-int.cpp (+116)
``````````diff
diff --git a/clang/include/clang/AST/OpenACCClause.h b/clang/include/clang/AST/OpenACCClause.h
index 081244fe0efb6..5f06117d65a47 100644
--- a/clang/include/clang/AST/OpenACCClause.h
+++ b/clang/include/clang/AST/OpenACCClause.h
@@ -842,9 +842,7 @@ struct OpenACCPrivateRecipe {
VarDecl *AllocaDecl;
Expr *InitExpr;
- OpenACCPrivateRecipe(VarDecl *A, Expr *I) : AllocaDecl(A), InitExpr(I) {
- assert(!AllocaDecl || AllocaDecl->getInit() == nullptr);
- }
+ OpenACCPrivateRecipe(VarDecl *A, Expr *I) : AllocaDecl(A), InitExpr(I) {}
bool isSet() const { return AllocaDecl; }
diff --git a/clang/lib/CIR/CodeGen/CIRGenFunction.h b/clang/lib/CIR/CodeGen/CIRGenFunction.h
index a0c571a544322..7413f5c8b2f79 100644
--- a/clang/lib/CIR/CodeGen/CIRGenFunction.h
+++ b/clang/lib/CIR/CodeGen/CIRGenFunction.h
@@ -1715,8 +1715,15 @@ class CIRGenFunction : public CIRGenTypeCache {
mlir::Location beginLoc;
mlir::Value varValue;
std::string name;
+ // The type of the original variable reference: that is, after 'bounds' have
+ // removed pointers/array types/etc. So in the case of int arr[5], and a
+ // private(arr[1]), 'origType' is 'int', but 'baseType' is 'int[5]'.
+ QualType origType;
QualType baseType;
llvm::SmallVector<mlir::Value> bounds;
+ // The list of types that we found when going through the bounds, which we
+ // can use to properly set the alloca section.
+ llvm::SmallVector<QualType> boundTypes;
};
// Gets the collection of info required to lower and OpenACC clause or cache
// construct variable reference.
diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenACC.cpp b/clang/lib/CIR/CodeGen/CIRGenOpenACC.cpp
index 907cb5fa11401..7f9350a9e4173 100644
--- a/clang/lib/CIR/CodeGen/CIRGenOpenACC.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenOpenACC.cpp
@@ -68,14 +68,33 @@ mlir::Value CIRGenFunction::createOpenACCConstantInt(mlir::Location loc,
CIRGenFunction::OpenACCDataOperandInfo
CIRGenFunction::getOpenACCDataOperandInfo(const Expr *e) {
const Expr *curVarExpr = e->IgnoreParenImpCasts();
+ QualType origType =
+ curVarExpr->getType().getNonReferenceType().getUnqualifiedType();
+ // Array sections are special, and we have to treat them that way.
+ if (const auto *section =
+ dyn_cast<ArraySectionExpr>(curVarExpr->IgnoreParenImpCasts()))
+ origType = ArraySectionExpr::getBaseOriginalType(section);
mlir::Location exprLoc = cgm.getLoc(curVarExpr->getBeginLoc());
llvm::SmallVector<mlir::Value> bounds;
+ llvm::SmallVector<QualType> boundTypes;
std::string exprString;
llvm::raw_string_ostream os(exprString);
e->printPretty(os, nullptr, getContext().getPrintingPolicy());
+ auto addBoundType = [&](const Expr *e) {
+ if (const auto *section = dyn_cast<ArraySectionExpr>(curVarExpr)) {
+ QualType baseTy = ArraySectionExpr::getBaseOriginalType(
+ section->getBase()->IgnoreParenImpCasts());
+ boundTypes.push_back(QualType(baseTy->getPointeeOrArrayElementType(), 0));
+ } else {
+ boundTypes.push_back(curVarExpr->getType());
+ }
+ };
+
+ addBoundType(curVarExpr);
+
while (isa<ArraySectionExpr, ArraySubscriptExpr>(curVarExpr)) {
mlir::Location boundLoc = cgm.getLoc(curVarExpr->getBeginLoc());
mlir::Value lowerBound;
@@ -115,19 +134,28 @@ CIRGenFunction::getOpenACCDataOperandInfo(const Expr *e) {
bounds.push_back(createBound(*this, this->builder, boundLoc, lowerBound,
upperBound, extent));
+ addBoundType(curVarExpr);
}
if (const auto *memExpr = dyn_cast<MemberExpr>(curVarExpr))
- return {exprLoc, emitMemberExpr(memExpr).getPointer(), exprString,
+ return {exprLoc,
+ emitMemberExpr(memExpr).getPointer(),
+ exprString,
+ origType,
curVarExpr->getType().getNonReferenceType().getUnqualifiedType(),
- std::move(bounds)};
+ std::move(bounds),
+ std::move(boundTypes)};
// Sema has made sure that only 4 types of things can get here, array
// subscript, array section, member expr, or DRE to a var decl (or the
// former 3 wrapping a var-decl), so we should be able to assume this is
// right.
const auto *dre = cast<DeclRefExpr>(curVarExpr);
- return {exprLoc, emitDeclRefLValue(dre).getPointer(), exprString,
+ return {exprLoc,
+ emitDeclRefLValue(dre).getPointer(),
+ exprString,
+ origType,
curVarExpr->getType().getNonReferenceType().getUnqualifiedType(),
- std::move(bounds)};
+ std::move(bounds),
+ std::move(boundTypes)};
}
diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
index 9959cf6c15792..dd37b101e9735 100644
--- a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
@@ -988,20 +988,16 @@ class OpenACCClauseCIREmitter final
{
mlir::OpBuilder::InsertionGuard guardCase(builder);
- // TODO: OpenACC: At the moment this is a bit of a hacky way of doing
- // this, and won't work when we get to bounds/etc. Do this for now to
- // limit the scope of this refactor.
- VarDecl *allocaDecl = varRecipe.AllocaDecl;
- allocaDecl->setInit(varRecipe.InitExpr);
- allocaDecl->setInitStyle(VarDecl::CallInit);
auto recipe =
OpenACCRecipeBuilder<mlir::acc::PrivateRecipeOp>(cgf, builder)
- .getOrCreateRecipe(cgf.getContext(), varExpr, allocaDecl,
- /*temporary=*/nullptr,
- OpenACCReductionOperator::Invalid,
- Decl::castToDeclContext(cgf.curFuncDecl),
- opInfo.baseType, privateOp.getResult());
+ .getOrCreateRecipe(
+ cgf.getContext(), varExpr, varRecipe.AllocaDecl,
+ varRecipe.InitExpr,
+ /*temporary=*/nullptr, OpenACCReductionOperator::Invalid,
+ Decl::castToDeclContext(cgf.curFuncDecl), opInfo.origType,
+ opInfo.bounds.size(), opInfo.boundTypes, opInfo.baseType,
+ privateOp.getResult());
// TODO: OpenACC: The dialect is going to change in the near future to
// have these be on a different operation, so when that changes, we
// probably need to change these here.
@@ -1042,12 +1038,13 @@ class OpenACCClauseCIREmitter final
auto recipe =
OpenACCRecipeBuilder<mlir::acc::FirstprivateRecipeOp>(cgf,
builder)
- .getOrCreateRecipe(cgf.getContext(), varExpr, allocaDecl,
- varRecipe.InitFromTemporary,
- OpenACCReductionOperator::Invalid,
- Decl::castToDeclContext(cgf.curFuncDecl),
- opInfo.baseType,
- firstPrivateOp.getResult());
+ .getOrCreateRecipe(
+ cgf.getContext(), varExpr, varRecipe.AllocaDecl,
+ varRecipe.InitExpr, varRecipe.InitFromTemporary,
+ OpenACCReductionOperator::Invalid,
+ Decl::castToDeclContext(cgf.curFuncDecl), opInfo.origType,
+ opInfo.bounds.size(), opInfo.boundTypes, opInfo.baseType,
+ firstPrivateOp.getResult());
// TODO: OpenACC: The dialect is going to change in the near future to
// have these be on a different operation, so when that changes, we
@@ -1089,11 +1086,13 @@ class OpenACCClauseCIREmitter final
auto recipe =
OpenACCRecipeBuilder<mlir::acc::ReductionRecipeOp>(cgf, builder)
- .getOrCreateRecipe(cgf.getContext(), varExpr, allocaDecl,
- /*temporary=*/nullptr,
- clause.getReductionOp(),
- Decl::castToDeclContext(cgf.curFuncDecl),
- opInfo.baseType, reductionOp.getResult());
+ .getOrCreateRecipe(
+ cgf.getContext(), varExpr, varRecipe.AllocaDecl,
+ varRecipe.InitExpr,
+ /*temporary=*/nullptr, clause.getReductionOp(),
+ Decl::castToDeclContext(cgf.curFuncDecl), opInfo.origType,
+ opInfo.bounds.size(), opInfo.boundTypes, opInfo.baseType,
+ reductionOp.getResult());
operation.addReduction(builder.getContext(), reductionOp, recipe);
}
diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenACCRecipe.h b/clang/lib/CIR/CodeGen/CIRGenOpenACCRecipe.h
index 102fd890e5579..4e60f892e7ab2 100644
--- a/clang/lib/CIR/CodeGen/CIRGenOpenACCRecipe.h
+++ b/clang/lib/CIR/CodeGen/CIRGenOpenACCRecipe.h
@@ -22,10 +22,140 @@
#include "mlir/Dialect/OpenACC/OpenACC.h"
namespace clang::CIRGen {
+
template <typename RecipeTy> class OpenACCRecipeBuilder {
CIRGen::CIRGenFunction &cgf;
CIRGen::CIRGenBuilderTy &builder;
+mlir::Block *createRecipeBlock(mlir::Region ®ion, mlir::Type opTy,
+ mlir::Location loc, size_t numBounds,
+ bool isInit) {
+ llvm::SmallVector<mlir::Type> types;
+ types.reserve(numBounds + 2);
+ types.push_back(opTy);
+ // The init section is the only one that doesn't have TWO copies of the
+ // operation-type. Copy has a to/from, and destroy has a
+ // 'reference'/'privatized' copy version.
+ if (!isInit)
+ types.push_back(opTy);
+
+ auto boundsTy = mlir::acc::DataBoundsType::get(&cgf.getMLIRContext());
+ for (size_t i = 0; i < numBounds; ++i)
+ types.push_back(boundsTy);
+
+ llvm::SmallVector<mlir::Location> locs{types.size(), loc};
+ return builder.createBlock(®ion, region.end(), types, locs);
+}
+// Creates a loop through an 'acc.bounds', leaving the 'insertion' point to be
+// the inside of the loop body. Traverses LB->UB UNLESS `inverse` is set.
+// Returns the 'subscriptedValue' changed with the new bounds subscript.
+mlir::Value createBoundsLoop(mlir::Value subscriptedValue, mlir::Value bound,
+ mlir::Location loc, bool inverse) {
+ mlir::Operation *bodyInsertLoc;
+
+ mlir::Type itrTy = cgf.cgm.convertType(cgf.getContext().UnsignedLongLongTy);
+ auto itrPtrTy = cir::PointerType::get(itrTy);
+ mlir::IntegerAttr itrAlign =
+ cgf.cgm.getSize(cgf.getContext().getTypeAlignInChars(
+ cgf.getContext().UnsignedLongLongTy));
+ auto idxType = mlir::IndexType::get(&cgf.getMLIRContext());
+
+ auto doSubscriptOp = [&](mlir::Value subVal,
+ cir::LoadOp idxLoad) -> mlir::Value {
+ auto eltTy = cast<cir::PointerType>(subVal.getType()).getPointee();
+
+ if (auto arrayTy = dyn_cast<cir::ArrayType>(eltTy))
+ return builder.getArrayElement(loc, loc, subVal, arrayTy.getElementType(),
+ idxLoad.getResult(),
+ /*shouldDecay=*/true);
+
+ assert(isa<cir::PointerType>(eltTy));
+
+ auto eltLoad = cir::LoadOp::create(builder, loc, {subVal});
+
+ return cir::PtrStrideOp::create(builder, loc, eltLoad.getType(), eltLoad,
+ idxLoad.getResult())
+ .getResult();
+ };
+
+ auto forStmtBuilder = [&]() {
+ // get the lower and upper bound for iterating over.
+ auto lowerBoundVal =
+ mlir::acc::GetLowerboundOp::create(builder, loc, idxType, bound);
+ auto lbConversion = mlir::UnrealizedConversionCastOp::create(
+ builder, loc, itrTy, lowerBoundVal.getResult());
+ auto upperBoundVal =
+ mlir::acc::GetUpperboundOp::create(builder, loc, idxType, bound);
+ auto ubConversion = mlir::UnrealizedConversionCastOp::create(
+ builder, loc, itrTy, upperBoundVal.getResult());
+
+ // Create a memory location for the iterator.
+ auto itr =
+ cir::AllocaOp::create(builder, loc, itrPtrTy, itrTy, "iter", itrAlign);
+ // Store to the iterator: either lower bound, or if inverse loop, upper
+ // bound.
+ if (inverse) {
+ cir::ConstantOp constOne = builder.getConstInt(loc, itrTy, 1);
+
+ auto sub =
+ cir::BinOp::create(builder, loc, itrTy, cir::BinOpKind::Sub,
+ ubConversion.getResult(0), constOne.getResult());
+
+ // Upperbound is exclusive, so subtract 1.
+ builder.CIRBaseBuilderTy::createStore(loc, sub.getResult(), itr);
+ } else {
+ // Lowerbound is inclusive, so we can include it.
+ builder.CIRBaseBuilderTy::createStore(loc, lbConversion.getResult(0),
+ itr);
+ }
+ // Save the 'end' iterator based on whether we are inverted or not. This
+ // end iterator never changes, so we can just get it and convert it, so no
+ // need to store/load/etc.
+ auto endItr = inverse ? lbConversion : ubConversion;
+
+ builder.createFor(
+ loc,
+ /*condBuilder=*/
+ [&](mlir::OpBuilder &b, mlir::Location loc) {
+ auto loadCur = cir::LoadOp::create(builder, loc, {itr});
+ // Use 'not equal' since we are just doing an increment/decrement.
+ auto cmp = builder.createCompare(
+ loc, inverse ? cir::CmpOpKind::ge : cir::CmpOpKind::lt,
+ loadCur.getResult(), endItr.getResult(0));
+ builder.createCondition(cmp);
+ },
+ /*bodyBuilder=*/
+ [&](mlir::OpBuilder &b, mlir::Location loc) {
+ auto load = cir::LoadOp::create(builder, loc, {itr});
+
+ if (subscriptedValue)
+ subscriptedValue = doSubscriptOp(subscriptedValue, load);
+ bodyInsertLoc = builder.createYield(loc);
+ },
+ /*stepBuilder=*/
+ [&](mlir::OpBuilder &b, mlir::Location loc) {
+ auto load = cir::LoadOp::create(builder, loc, {itr});
+ auto unary = cir::UnaryOp::create(builder, loc, load.getType(),
+ inverse ? cir::UnaryOpKind::Dec
+ : cir::UnaryOpKind::Inc,
+ load.getResult());
+ builder.CIRBaseBuilderTy::createStore(loc, unary.getResult(), itr);
+ builder.createYield(loc);
+ });
+ };
+
+ cir::ScopeOp::create(builder, loc,
+ [&](mlir::OpBuilder &b, mlir::Location loc) {
+ forStmtBuilder();
+ builder.createYield(loc);
+ });
+
+ // Leave the insertion point to be inside the body, so we can loop over
+ // these things.
+ builder.setInsertionPoint(bodyInsertLoc);
+ return subscriptedValue;
+}
+
mlir::acc::ReductionOperator convertReductionOp(OpenACCReductionOperator op) {
switch (op) {
case OpenACCReductionOperator::Addition:
@@ -54,6 +184,7 @@ template <typename RecipeTy> class OpenACCRecipeBuilder {
}
std::string getRecipeName(SourceRange loc, QualType baseType,
+ unsigned numBounds,
OpenACCReductionOperator reductionOp) {
std::string recipeName;
{
@@ -106,6 +237,11 @@ template <typename RecipeTy> class OpenACCRecipeBuilder {
static_assert(!sizeof(RecipeTy), "Unknown Recipe op kind");
}
+ // The naming convention from Flang with bounds doesn't map to C++ types
+ // very well, so we're just going to choose our own here.
+ if (numBounds)
+ stream << "_Bcnt" << numBounds << '_';
+
MangleContext &mc = cgf.cgm.getCXXABI().getMangleContext();
mc.mangleCanonicalTypeName(baseType, stream);
}
@@ -117,9 +253,9 @@ template <typename RecipeTy> class OpenACCRecipeBuilder {
CIRGenFunction::AutoVarEmission tempDeclEmission,
mlir::acc::FirstprivateRecipeOp recipe, const VarDecl *varRecipe,
const VarDecl *temporary) {
- mlir::Block *block = builder.createBlock(
- &recipe.getCopyRegion(), recipe.getCopyRegion().end(),
- {mainOp.getType(), mainOp.getType()}, {loc, loc});
+ mlir::Block *block =
+ createRecipeBlock(recipe.getCopyRegion(), mainOp.getType(), loc,
+ /*numBounds=*/0, /*isInit=*/false);
builder.setInsertionPointToEnd(&recipe.getCopyRegion().back());
CIRGenFunction::LexicalScope ls(cgf, loc, block);
@@ -143,6 +279,54 @@ template <typename RecipeTy> class OpenACCRecipeBuilder {
cgf.emitAutoVarInit(tempDeclEmission);
mlir::acc::YieldOp::create(builder, locEnd);
}
+ // TODO: OpenACC: When we get this implemented for the reduction/firstprivate,
+ // this might end up re-merging with createRecipeInitCopy. For now, keep it
+ // separate until we're sure what everything looks like to keep this as clean
+ // as possible.
+ void createPrivateInitRecipe(mlir::Location loc, mlir::Location locEnd,
+ SourceRange exprRange, mlir::Value mainOp,
+ mlir::acc::PrivateRecipeOp recipe,
+ size_t numBounds,
+ llvm::ArrayRef<QualType> boundTypes,
+ const VarDecl *allocaDecl, QualType origType,
+ const Expr *initExpr) {
+ assert(allocaDecl && "Required recipe variable not set?");
+ CIRGenFunction::DeclMapRevertingRAII declMapRAII{cgf, allocaDecl};
+
+ mlir::Block *block =
+ createRecipeBlock(recipe.getInitRegion(), mainOp.getType(), loc,
+ numBounds, /*isInit=*/true);
+ builder.setInsertionPointToEnd(&recipe.getInitRegion().back());
+ CIRGenFunction::LexicalScope ls(cgf, loc, block);
+
+ const Type *allocaPointeeType =
+ allocaDecl->getType()->getPointeeOrArrayElementType();
+ // We are OK with no init for builtins, arrays of builtins, or pointers,
+ // else we should NYI so we know to go look for these.
+ if (cgf.getContext().getLangOpts().CPlusPlus && !allocaDecl->getInit() &&
+ !allocaDecl->getType()->isPointerType() &&
+ !allocaPointeeType->isBuiltinType() &&
+ !allocaPointeeType->isPointerType()) {
+ // If we don't have any initialization recipe, we failed during Sema to
+ // initialize this correctly. If we disable the
+ // Sema::TentativeAnalysisScopes in SemaOpenACC::CreateInitRecipe, it'll
+ // emit an error to tell us. However, emitting those errors during
+ // production is a violation of the standard, so we cannot do them.
+ cgf.cgm.errorNYI(exprRange, "private default-init recipe");
+ }
+
+ if (!numBounds) {
+ // This is an 'easy' case, we just have to use the builtin init stuff to
+ // initialize this variable correctly.
+ CIRGenFunction::AutoVarEmission tempDeclEmission =
+ cgf.emitAutoVarAlloca(*allocaDecl, builder.saveInsertionPoint());
+ cgf.emitAutoVarInit(tempDeclEmission);
+ } else {
+ cgf.cgm.errorNYI(exprRange, "private-init with bounds");
+ }
+
+ mlir::acc::YieldOp::create(builder, locEnd);
+ }
// Create the 'init' section of the recipe, including the 'copy' section for
// 'firstprivate'. Note that this function is not 'insertion point' clean, in
@@ -160,9 +344,9 @@ template <typename RecipeTy> class OpenACCRecipeBuilder {
// Do the 'init' section of the recipe IR, which does an alloca, then the
// initialization (except for firstprivate).
- mlir::Block *block = builder.createBlock(&recipe.getInitRegion(),
- recipe.getInitRegion().end(),
- {mainOp.getType()}, {loc});
+ mlir::Block *block =
+ createRecipeBlock(recipe.getInitRegion(), mainOp.getType(), loc,
+ /*numBounds=*/0, /*isInit=*/true);
builder.setInsertionPointToEnd(&recipe.getInitRegion().back());
CIRGenFunction::LexicalScope ls(cgf, loc, block);
@@ -241,22 +425,42 @@ template <typename RecipeTy> class OpenACCRecipeBuilder {
// doesn't restore it aftewards.
void createRecipeDestroySection(mlir::Location loc, mlir::Location locEnd,
...
[truncated]
``````````
</details>
https://github.com/llvm/llvm-project/pull/160189
More information about the cfe-commits
mailing list