[clang] 66eadbb - [OpenACC][CIR] Implement 'init' lowering for private clause vars (#151781)
via cfe-commits
cfe-commits at lists.llvm.org
Mon Aug 4 11:15:01 PDT 2025
Author: Erich Keane
Date: 2025-08-04T11:14:58-07:00
New Revision: 66eadbb235320b476214c810f5bbdc0daa44d2d7
URL: https://github.com/llvm/llvm-project/commit/66eadbb235320b476214c810f5bbdc0daa44d2d7
DIFF: https://github.com/llvm/llvm-project/commit/66eadbb235320b476214c810f5bbdc0daa44d2d7.diff
LOG: [OpenACC][CIR] Implement 'init' lowering for private clause vars (#151781)
Previously, #151360 implemented 'private' clause lowering, but didn't
properly initialize the variables. This patch adds that behavior to make
sure we correctly get the constructor or other init called.
Added:
clang/test/CIR/CodeGenOpenACC/compute-private-clause-templates.cpp
Modified:
clang/include/clang/AST/OpenACCClause.h
clang/include/clang/Sema/SemaOpenACC.h
clang/lib/AST/OpenACCClause.cpp
clang/lib/AST/StmtProfile.cpp
clang/lib/CIR/CodeGen/CIRGenOpenACC.cpp
clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
clang/lib/Sema/SemaOpenACC.cpp
clang/lib/Sema/SemaOpenACCClause.cpp
clang/lib/Sema/TreeTransform.h
clang/lib/Serialization/ASTReader.cpp
clang/lib/Serialization/ASTWriter.cpp
clang/test/CIR/CodeGenOpenACC/combined-private-clause.cpp
clang/test/CIR/CodeGenOpenACC/compute-private-clause.cpp
clang/test/CIR/CodeGenOpenACC/loop-private-clause.cpp
clang/tools/libclang/CIndex.cpp
Removed:
################################################################################
diff --git a/clang/include/clang/AST/OpenACCClause.h b/clang/include/clang/AST/OpenACCClause.h
index 71ad24a427105..32a06bc84ad20 100644
--- a/clang/include/clang/AST/OpenACCClause.h
+++ b/clang/include/clang/AST/OpenACCClause.h
@@ -837,23 +837,43 @@ class OpenACCClauseWithVarList : public OpenACCClauseWithExprs {
class OpenACCPrivateClause final
: public OpenACCClauseWithVarList,
- private llvm::TrailingObjects<OpenACCPrivateClause, Expr *> {
+ private llvm::TrailingObjects<OpenACCPrivateClause, Expr *, VarDecl *> {
friend TrailingObjects;
OpenACCPrivateClause(SourceLocation BeginLoc, SourceLocation LParenLoc,
- ArrayRef<Expr *> VarList, SourceLocation EndLoc)
+ ArrayRef<Expr *> VarList,
+ ArrayRef<VarDecl *> InitRecipes, SourceLocation EndLoc)
: OpenACCClauseWithVarList(OpenACCClauseKind::Private, BeginLoc,
LParenLoc, EndLoc) {
- setExprs(getTrailingObjects(VarList.size()), VarList);
+ assert(VarList.size() == InitRecipes.size());
+ setExprs(getTrailingObjects<Expr *>(VarList.size()), VarList);
+ llvm::uninitialized_copy(InitRecipes, getTrailingObjects<VarDecl *>());
}
public:
static bool classof(const OpenACCClause *C) {
return C->getClauseKind() == OpenACCClauseKind::Private;
}
+ // Gets a list of 'made up' `VarDecl` objects that can be used by codegen to
+ // ensure that we properly initialize each of these variables.
+ ArrayRef<VarDecl *> getInitRecipes() {
+ return ArrayRef<VarDecl *>{getTrailingObjects<VarDecl *>(),
+ getExprs().size()};
+ }
+
+ ArrayRef<VarDecl *> getInitRecipes() const {
+ return ArrayRef<VarDecl *>{getTrailingObjects<VarDecl *>(),
+ getExprs().size()};
+ }
+
static OpenACCPrivateClause *
Create(const ASTContext &C, SourceLocation BeginLoc, SourceLocation LParenLoc,
- ArrayRef<Expr *> VarList, SourceLocation EndLoc);
+ ArrayRef<Expr *> VarList, ArrayRef<VarDecl *> InitRecipes,
+ SourceLocation EndLoc);
+
+ size_t numTrailingObjects(OverloadToken<Expr *>) const {
+ return getExprs().size();
+ }
};
class OpenACCFirstPrivateClause final
diff --git a/clang/include/clang/Sema/SemaOpenACC.h b/clang/include/clang/Sema/SemaOpenACC.h
index b7e7f5d97bcef..f51045d26e23b 100644
--- a/clang/include/clang/Sema/SemaOpenACC.h
+++ b/clang/include/clang/Sema/SemaOpenACC.h
@@ -241,6 +241,10 @@ class SemaOpenACC : public SemaBase {
SourceLocation ClauseLoc,
ArrayRef<const OpenACCClause *> Clauses);
+ // Creates a VarDecl with a proper default init for the purposes of a
+ // `private` clause, so it can be used to generate a recipe later.
+ VarDecl *CreateInitRecipe(const Expr *VarExpr);
+
public:
ComputeConstructInfo &getActiveComputeConstructInfo() {
return ActiveComputeConstructInfo;
diff --git a/clang/lib/AST/OpenACCClause.cpp b/clang/lib/AST/OpenACCClause.cpp
index 60ec10a986e5e..f21e645697656 100644
--- a/clang/lib/AST/OpenACCClause.cpp
+++ b/clang/lib/AST/OpenACCClause.cpp
@@ -314,14 +314,17 @@ OpenACCTileClause *OpenACCTileClause::Create(const ASTContext &C,
return new (Mem) OpenACCTileClause(BeginLoc, LParenLoc, SizeExprs, EndLoc);
}
-OpenACCPrivateClause *OpenACCPrivateClause::Create(const ASTContext &C,
- SourceLocation BeginLoc,
- SourceLocation LParenLoc,
- ArrayRef<Expr *> VarList,
- SourceLocation EndLoc) {
- void *Mem = C.Allocate(
- OpenACCPrivateClause::totalSizeToAlloc<Expr *>(VarList.size()));
- return new (Mem) OpenACCPrivateClause(BeginLoc, LParenLoc, VarList, EndLoc);
+OpenACCPrivateClause *
+OpenACCPrivateClause::Create(const ASTContext &C, SourceLocation BeginLoc,
+ SourceLocation LParenLoc, ArrayRef<Expr *> VarList,
+ ArrayRef<VarDecl *> InitRecipes,
+ SourceLocation EndLoc) {
+ assert(VarList.size() == InitRecipes.size());
+ void *Mem =
+ C.Allocate(OpenACCPrivateClause::totalSizeToAlloc<Expr *, VarDecl *>(
+ VarList.size(), InitRecipes.size()));
+ return new (Mem)
+ OpenACCPrivateClause(BeginLoc, LParenLoc, VarList, InitRecipes, EndLoc);
}
OpenACCFirstPrivateClause *OpenACCFirstPrivateClause::Create(
diff --git a/clang/lib/AST/StmtProfile.cpp b/clang/lib/AST/StmtProfile.cpp
index c61450e19f1b6..57834ca4a89ca 100644
--- a/clang/lib/AST/StmtProfile.cpp
+++ b/clang/lib/AST/StmtProfile.cpp
@@ -2636,6 +2636,9 @@ void OpenACCClauseProfiler::VisitCollapseClause(
void OpenACCClauseProfiler::VisitPrivateClause(
const OpenACCPrivateClause &Clause) {
VisitClauseWithVarList(Clause);
+
+ for (auto *VD : Clause.getInitRecipes())
+ Profiler.VisitDecl(VD);
}
void OpenACCClauseProfiler::VisitFirstPrivateClause(
diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenACC.cpp b/clang/lib/CIR/CodeGen/CIRGenOpenACC.cpp
index 32095cb687e88..907cb5fa11401 100644
--- a/clang/lib/CIR/CodeGen/CIRGenOpenACC.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenOpenACC.cpp
@@ -119,7 +119,8 @@ CIRGenFunction::getOpenACCDataOperandInfo(const Expr *e) {
if (const auto *memExpr = dyn_cast<MemberExpr>(curVarExpr))
return {exprLoc, emitMemberExpr(memExpr).getPointer(), exprString,
- curVarExpr->getType(), std::move(bounds)};
+ curVarExpr->getType().getNonReferenceType().getUnqualifiedType(),
+ std::move(bounds)};
// 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
@@ -127,5 +128,6 @@ CIRGenFunction::getOpenACCDataOperandInfo(const Expr *e) {
// right.
const auto *dre = cast<DeclRefExpr>(curVarExpr);
return {exprLoc, emitDeclRefLValue(dre).getPointer(), exprString,
- curVarExpr->getType(), std::move(bounds)};
+ curVarExpr->getType().getNonReferenceType().getUnqualifiedType(),
+ std::move(bounds)};
}
diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
index 5a6e66550c0bd..bb9054a68b5c7 100644
--- a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
@@ -358,8 +358,8 @@ class OpenACCClauseCIREmitter final
template <typename RecipeTy>
RecipeTy getOrCreateRecipe(ASTContext &astCtx, const Expr *varRef,
- DeclContext *dc, QualType baseType,
- mlir::Value mainOp) {
+ const VarDecl *varRecipe, DeclContext *dc,
+ QualType baseType, mlir::Value mainOp) {
mlir::ModuleOp mod =
builder.getBlock()->getParent()->getParentOfType<mlir::ModuleOp>();
@@ -398,12 +398,6 @@ class OpenACCClauseCIREmitter final
auto recipe =
RecipeTy::create(modBuilder, loc, recipeName, mainOp.getType());
- // Magic-up a var-decl so we can use normal init/destruction operations for
- // a variable declaration.
- VarDecl &tempDecl = *VarDecl::Create(
- astCtx, dc, varRef->getBeginLoc(), varRef->getBeginLoc(),
- &astCtx.Idents.get("openacc.private.init"), baseType,
- astCtx.getTrivialTypeSourceInfo(baseType), SC_Auto);
CIRGenFunction::AutoVarEmission tempDeclEmission{
CIRGenFunction::AutoVarEmission::invalid()};
@@ -422,9 +416,11 @@ class OpenACCClauseCIREmitter final
"OpenACC non-private recipe init");
}
- tempDeclEmission =
- cgf.emitAutoVarAlloca(tempDecl, builder.saveInsertionPoint());
- cgf.emitAutoVarInit(tempDeclEmission);
+ if (varRecipe) {
+ tempDeclEmission =
+ cgf.emitAutoVarAlloca(*varRecipe, builder.saveInsertionPoint());
+ cgf.emitAutoVarInit(tempDeclEmission);
+ }
mlir::acc::YieldOp::create(builder, locEnd);
}
@@ -439,7 +435,7 @@ class OpenACCClauseCIREmitter final
}
// Destroy section (doesn't currently exist).
- if (tempDecl.needsDestruction(cgf.getContext())) {
+ if (varRecipe && varRecipe->needsDestruction(cgf.getContext())) {
llvm::SmallVector<mlir::Type> argsTys{mainOp.getType()};
llvm::SmallVector<mlir::Location> argsLocs{loc};
mlir::Block *block = builder.createBlock(&recipe.getDestroyRegion(),
@@ -450,7 +446,7 @@ class OpenACCClauseCIREmitter final
mlir::Type elementTy =
mlir::cast<cir::PointerType>(mainOp.getType()).getPointee();
Address addr{block->getArgument(0), elementTy,
- cgf.getContext().getDeclAlign(&tempDecl)};
+ cgf.getContext().getDeclAlign(varRecipe)};
cgf.emitDestroy(addr, baseType,
cgf.getDestroyer(QualType::DK_cxx_destructor));
@@ -1080,9 +1076,10 @@ class OpenACCClauseCIREmitter final
void VisitPrivateClause(const OpenACCPrivateClause &clause) {
if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
mlir::acc::LoopOp>) {
- for (const Expr *var : clause.getVarList()) {
+ for (const auto [varExpr, varRecipe] :
+ llvm::zip_equal(clause.getVarList(), clause.getInitRecipes())) {
CIRGenFunction::OpenACCDataOperandInfo opInfo =
- cgf.getOpenACCDataOperandInfo(var);
+ cgf.getOpenACCDataOperandInfo(varExpr);
auto privateOp = mlir::acc::PrivateOp::create(
builder, opInfo.beginLoc, opInfo.varValue, /*structured=*/true,
/*implicit=*/false, opInfo.name, opInfo.bounds);
@@ -1091,8 +1088,9 @@ class OpenACCClauseCIREmitter final
{
mlir::OpBuilder::InsertionGuard guardCase(builder);
auto recipe = getOrCreateRecipe<mlir::acc::PrivateRecipeOp>(
- cgf.getContext(), var, Decl::castToDeclContext(cgf.curFuncDecl),
- opInfo.baseType, privateOp.getResult());
+ cgf.getContext(), varExpr, varRecipe,
+ Decl::castToDeclContext(cgf.curFuncDecl), opInfo.baseType,
+ privateOp.getResult());
// TODO: OpenACC: The dialect is going to change in the near future to
// have these be on a
diff erent operation, so when that changes, we
// probably need to change these here.
diff --git a/clang/lib/Sema/SemaOpenACC.cpp b/clang/lib/Sema/SemaOpenACC.cpp
index 8bfea623103e4..8f32817aec48f 100644
--- a/clang/lib/Sema/SemaOpenACC.cpp
+++ b/clang/lib/Sema/SemaOpenACC.cpp
@@ -17,6 +17,7 @@
#include "clang/Basic/DiagnosticSema.h"
#include "clang/Basic/OpenACCKinds.h"
#include "clang/Basic/SourceManager.h"
+#include "clang/Sema/Initialization.h"
#include "clang/Sema/Scope.h"
#include "clang/Sema/Sema.h"
#include "llvm/ADT/StringExtras.h"
@@ -2552,3 +2553,50 @@ ExprResult
SemaOpenACC::ActOnOpenACCAsteriskSizeExpr(SourceLocation AsteriskLoc) {
return BuildOpenACCAsteriskSizeExpr(AsteriskLoc);
}
+
+VarDecl *SemaOpenACC::CreateInitRecipe(const Expr *VarExpr) {
+ // Strip off any array subscripts/array section exprs to get to the type of
+ // the variable.
+ while (isa_and_present<ArraySectionExpr, ArraySubscriptExpr>(VarExpr)) {
+ if (const auto *AS = dyn_cast<ArraySectionExpr>(VarExpr))
+ VarExpr = AS->getBase()->IgnoreParenImpCasts();
+ else if (const auto *Sub = dyn_cast<ArraySubscriptExpr>(VarExpr))
+ VarExpr = Sub->getBase()->IgnoreParenImpCasts();
+ }
+
+ // If for some reason the expression is invalid, or this is dependent, just
+ // fill in with nullptr. We'll count on TreeTransform to make this if
+ // necessary.
+ if (!VarExpr || VarExpr->getType()->isDependentType())
+ return nullptr;
+
+ QualType VarTy =
+ VarExpr->getType().getNonReferenceType().getUnqualifiedType();
+
+ VarDecl *Recipe = VarDecl::Create(
+ getASTContext(), SemaRef.getCurContext(), VarExpr->getBeginLoc(),
+ VarExpr->getBeginLoc(),
+ &getASTContext().Idents.get("openacc.private.init"), VarTy,
+ getASTContext().getTrivialTypeSourceInfo(VarTy), SC_Auto);
+
+ ExprResult Init;
+
+ {
+ // Trap errors so we don't get weird ones here. If we can't init, we'll just
+ // swallow the errors.
+ Sema::TentativeAnalysisScope Trap{SemaRef};
+ InitializedEntity Entity = InitializedEntity::InitializeVariable(Recipe);
+ InitializationKind Kind =
+ InitializationKind::CreateDefault(Recipe->getLocation());
+
+ InitializationSequence InitSeq(SemaRef.SemaRef, Entity, Kind, {});
+ Init = InitSeq.Perform(SemaRef.SemaRef, Entity, Kind, {});
+ }
+
+ if (Init.get()) {
+ Recipe->setInit(Init.get());
+ Recipe->setInitStyle(VarDecl::CallInit);
+ }
+
+ return Recipe;
+}
diff --git a/clang/lib/Sema/SemaOpenACCClause.cpp b/clang/lib/Sema/SemaOpenACCClause.cpp
index b54a0124e9495..9f1a6177b945c 100644
--- a/clang/lib/Sema/SemaOpenACCClause.cpp
+++ b/clang/lib/Sema/SemaOpenACCClause.cpp
@@ -795,9 +795,15 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitPrivateClause(
// really isn't anything to do here. GCC does some duplicate-finding, though
// it isn't apparent in the standard where this is justified.
- return OpenACCPrivateClause::Create(Ctx, Clause.getBeginLoc(),
- Clause.getLParenLoc(),
- Clause.getVarList(), Clause.getEndLoc());
+ llvm::SmallVector<VarDecl *> InitRecipes;
+
+ // Assemble the recipes list.
+ for (const Expr *VarExpr : Clause.getVarList())
+ InitRecipes.push_back(SemaRef.CreateInitRecipe(VarExpr));
+
+ return OpenACCPrivateClause::Create(
+ Ctx, Clause.getBeginLoc(), Clause.getLParenLoc(), Clause.getVarList(),
+ InitRecipes, Clause.getEndLoc());
}
OpenACCClause *SemaOpenACCClauseVisitor::VisitFirstPrivateClause(
diff --git a/clang/lib/Sema/TreeTransform.h b/clang/lib/Sema/TreeTransform.h
index f2e800ad17dd3..5ff5fbf52ae25 100644
--- a/clang/lib/Sema/TreeTransform.h
+++ b/clang/lib/Sema/TreeTransform.h
@@ -11740,20 +11740,26 @@ class OpenACCClauseTransform final
SemaOpenACC::OpenACCParsedClause &ParsedClause;
OpenACCClause *NewClause = nullptr;
+ ExprResult VisitVar(Expr *VarRef) {
+ ExprResult Res = Self.TransformExpr(VarRef);
+
+ if (!Res.isUsable())
+ return Res;
+
+ Res = Self.getSema().OpenACC().ActOnVar(ParsedClause.getDirectiveKind(),
+ ParsedClause.getClauseKind(),
+ Res.get());
+
+ return Res;
+ }
+
llvm::SmallVector<Expr *> VisitVarList(ArrayRef<Expr *> VarList) {
llvm::SmallVector<Expr *> InstantiatedVarList;
for (Expr *CurVar : VarList) {
- ExprResult Res = Self.TransformExpr(CurVar);
-
- if (!Res.isUsable())
- continue;
+ ExprResult VarRef = VisitVar(CurVar);
- Res = Self.getSema().OpenACC().ActOnVar(ParsedClause.getDirectiveKind(),
- ParsedClause.getClauseKind(),
- Res.get());
-
- if (Res.isUsable())
- InstantiatedVarList.push_back(Res.get());
+ if (VarRef.isUsable())
+ InstantiatedVarList.push_back(VarRef.get());
}
return InstantiatedVarList;
@@ -11880,12 +11886,31 @@ void OpenACCClauseTransform<Derived>::VisitNumGangsClause(
template <typename Derived>
void OpenACCClauseTransform<Derived>::VisitPrivateClause(
const OpenACCPrivateClause &C) {
- ParsedClause.setVarListDetails(VisitVarList(C.getVarList()),
+ llvm::SmallVector<Expr *> InstantiatedVarList;
+ llvm::SmallVector<VarDecl *> InitRecipes;
+
+ for (const auto [RefExpr, InitRecipe] :
+ llvm::zip(C.getVarList(), C.getInitRecipes())) {
+ ExprResult VarRef = VisitVar(RefExpr);
+
+ if (VarRef.isUsable()) {
+ InstantiatedVarList.push_back(VarRef.get());
+
+ // We only have to create a new one if it is dependent, and Sema won't
+ // make one of these unless the type is non-dependent.
+ if (InitRecipe)
+ InitRecipes.push_back(InitRecipe);
+ else
+ InitRecipes.push_back(
+ Self.getSema().OpenACC().CreateInitRecipe(VarRef.get()));
+ }
+ }
+ ParsedClause.setVarListDetails(InstantiatedVarList,
OpenACCModifierKind::Invalid);
NewClause = OpenACCPrivateClause::Create(
Self.getSema().getASTContext(), ParsedClause.getBeginLoc(),
- ParsedClause.getLParenLoc(), ParsedClause.getVarList(),
+ ParsedClause.getLParenLoc(), ParsedClause.getVarList(), InitRecipes,
ParsedClause.getEndLoc());
}
diff --git a/clang/lib/Serialization/ASTReader.cpp b/clang/lib/Serialization/ASTReader.cpp
index 1b2aaad69eabd..95e92cf6790d7 100644
--- a/clang/lib/Serialization/ASTReader.cpp
+++ b/clang/lib/Serialization/ASTReader.cpp
@@ -12854,8 +12854,13 @@ OpenACCClause *ASTRecordReader::readOpenACCClause() {
case OpenACCClauseKind::Private: {
SourceLocation LParenLoc = readSourceLocation();
llvm::SmallVector<Expr *> VarList = readOpenACCVarList();
+
+ llvm::SmallVector<VarDecl *> RecipeList;
+ for (unsigned I = 0; I < VarList.size(); ++I)
+ RecipeList.push_back(readDeclAs<VarDecl>());
+
return OpenACCPrivateClause::Create(getContext(), BeginLoc, LParenLoc,
- VarList, EndLoc);
+ VarList, RecipeList, EndLoc);
}
case OpenACCClauseKind::Host: {
SourceLocation LParenLoc = readSourceLocation();
diff --git a/clang/lib/Serialization/ASTWriter.cpp b/clang/lib/Serialization/ASTWriter.cpp
index 2b40be3b7349d..955f7b914f5db 100644
--- a/clang/lib/Serialization/ASTWriter.cpp
+++ b/clang/lib/Serialization/ASTWriter.cpp
@@ -8740,6 +8740,9 @@ void ASTRecordWriter::writeOpenACCClause(const OpenACCClause *C) {
const auto *PC = cast<OpenACCPrivateClause>(C);
writeSourceLocation(PC->getLParenLoc());
writeOpenACCVarList(PC);
+
+ for (VarDecl *VD : PC->getInitRecipes())
+ AddDeclRef(VD);
return;
}
case OpenACCClauseKind::Host: {
diff --git a/clang/test/CIR/CodeGenOpenACC/combined-private-clause.cpp b/clang/test/CIR/CodeGenOpenACC/combined-private-clause.cpp
index 3306c55b4a8e7..0b58bae9239b2 100644
--- a/clang/test/CIR/CodeGenOpenACC/combined-private-clause.cpp
+++ b/clang/test/CIR/CodeGenOpenACC/combined-private-clause.cpp
@@ -43,7 +43,24 @@ struct HasDtor {
//
// CHECK-NEXT: acc.private.recipe @privatization__ZTSA5_14NonDefaultCtor : !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>> init {
// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>> {{.*}}):
-// CHECK-NEXT: cir.alloca !cir.array<!rec_NonDefaultCtor x 5>, !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>>, ["openacc.private.init"]
+// CHECK-NEXT: %[[ALLOCA:.*]] = cir.alloca !cir.array<!rec_NonDefaultCtor x 5>, !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>>, ["openacc.private.init", init]
+// CHECK-NEXT: %[[LAST_IDX:.*]] = cir.const #cir.int<5> : !u64i
+// CHECK-NEXT: %[[ARRPTR:.*]] = cir.cast(array_to_ptrdecay, %[[ALLOCA]] : !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>>), !cir.ptr<!rec_NonDefaultCtor>
+// CHECK-NEXT: %[[LAST_ELEM:.*]] = cir.ptr_stride(%[[ARRPTR]] : !cir.ptr<!rec_NonDefaultCtor>, %[[LAST_IDX]] : !u64i), !cir.ptr<!rec_NonDefaultCtor>
+// CHECK-NEXT: %[[ITR:.*]] = cir.alloca !cir.ptr<!rec_NonDefaultCtor>, !cir.ptr<!cir.ptr<!rec_NonDefaultCtor>>, ["__array_idx"]
+// CHECK-NEXT: cir.store %[[ARRPTR]], %[[ITR]] : !cir.ptr<!rec_NonDefaultCtor>, !cir.ptr<!cir.ptr<!rec_NonDefaultCtor>>
+// CHECK-NEXT: cir.do {
+// CHECK-NEXT: %[[ELEM_LOAD:.*]] = cir.load %[[ITR]] : !cir.ptr<!cir.ptr<!rec_NonDefaultCtor>>, !cir.ptr<!rec_NonDefaultCtor>
+// CHECK-NEXT: cir.call @_ZN14NonDefaultCtorC1Ev(%[[ELEM_LOAD]]) : (!cir.ptr<!rec_NonDefaultCtor>) -> ()
+// CHECK-NEXT: %[[ONE_CONST:.*]] = cir.const #cir.int<1> : !u64i
+// CHECK-NEXT: %[[ELEM:.*]] = cir.ptr_stride(%[[ELEM_LOAD]] : !cir.ptr<!rec_NonDefaultCtor>, %[[ONE_CONST]] : !u64i), !cir.ptr<!rec_NonDefaultCtor>
+// CHECK-NEXT: cir.store %[[ELEM]], %[[ITR]] : !cir.ptr<!rec_NonDefaultCtor>, !cir.ptr<!cir.ptr<!rec_NonDefaultCtor>>
+// CHECK-NEXT: cir.yield
+// CHECK-NEXT: } while {
+// CHECK-NEXT: %[[ELEM_LOAD:.*]] = cir.load %[[ITR]] : !cir.ptr<!cir.ptr<!rec_NonDefaultCtor>>, !cir.ptr<!rec_NonDefaultCtor>
+// CHECK-NEXT: %[[CMP:.*]] = cir.cmp(ne, %[[ELEM_LOAD]], %[[LAST_ELEM]]) : !cir.ptr<!rec_NonDefaultCtor>, !cir.bool
+// CHECK-NEXT: cir.condition(%[[CMP]])
+// CHECK-NEXT: }
// CHECK-NEXT: acc.yield
// CHECK-NEXT: }
//
@@ -83,7 +100,8 @@ struct HasDtor {
//
// CHECK-NEXT: acc.private.recipe @privatization__ZTS14NonDefaultCtor : !cir.ptr<!rec_NonDefaultCtor> init {
// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!rec_NonDefaultCtor> {{.*}}):
-// CHECK-NEXT: cir.alloca !rec_NonDefaultCtor, !cir.ptr<!rec_NonDefaultCtor>, ["openacc.private.init"]
+// CHECK-NEXT: %[[ALLOCA:.*]] = cir.alloca !rec_NonDefaultCtor, !cir.ptr<!rec_NonDefaultCtor>, ["openacc.private.init", init]
+// CHECK-NEXT: cir.call @_ZN14NonDefaultCtorC1Ev(%[[ALLOCA]]) : (!cir.ptr<!rec_NonDefaultCtor>) -> ()
// CHECK-NEXT: acc.yield
// CHECK-NEXT: }
//
diff --git a/clang/test/CIR/CodeGenOpenACC/compute-private-clause-templates.cpp b/clang/test/CIR/CodeGenOpenACC/compute-private-clause-templates.cpp
new file mode 100644
index 0000000000000..659e84bc892b6
--- /dev/null
+++ b/clang/test/CIR/CodeGenOpenACC/compute-private-clause-templates.cpp
@@ -0,0 +1,79 @@
+// RUN: %clang_cc1 -fopenacc -triple x86_64-linux-gnu -Wno-openacc-self-if-potential-conflict -emit-cir -fclangir -triple x86_64-linux-pc %s -o - | FileCheck %s
+
+struct CopyConstruct {
+ CopyConstruct() = default;
+ CopyConstruct(const CopyConstruct&);
+};
+
+struct NonDefaultCtor {
+ NonDefaultCtor();
+};
+
+struct HasDtor {
+ ~HasDtor();
+};
+
+// CHECK: acc.private.recipe @privatization__ZTSi : !cir.ptr<!s32i> init {
+// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!s32i> {{.*}}):
+// CHECK-NEXT: cir.alloca !s32i, !cir.ptr<!s32i>, ["openacc.private.init"]
+// CHECK-NEXT: acc.yield
+// CHECK-NEXT: }
+//
+// CHECK-NEXT: acc.private.recipe @privatization__ZTS7HasDtor : !cir.ptr<!rec_HasDtor> init {
+// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!rec_HasDtor> {{.*}}):
+// CHECK-NEXT: cir.alloca !rec_HasDtor, !cir.ptr<!rec_HasDtor>, ["openacc.private.init"]
+// CHECK-NEXT: acc.yield
+// CHECK-NEXT: } destroy {
+// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!rec_HasDtor> {{.*}}):
+// CHECK-NEXT: cir.call @_ZN7HasDtorD1Ev(%[[ARG]]) nothrow : (!cir.ptr<!rec_HasDtor>) -> ()
+// CHECK-NEXT: acc.yield
+// CHECK-NEXT: }
+//
+// CHECK-NEXT: acc.private.recipe @privatization__ZTS14NonDefaultCtor : !cir.ptr<!rec_NonDefaultCtor> init {
+// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!rec_NonDefaultCtor> {{.*}}):
+// CHECK-NEXT: %[[ALLOCA:.*]] = cir.alloca !rec_NonDefaultCtor, !cir.ptr<!rec_NonDefaultCtor>, ["openacc.private.init", init]
+// CHECK-NEXT: cir.call @_ZN14NonDefaultCtorC1Ev(%[[ALLOCA]]) : (!cir.ptr<!rec_NonDefaultCtor>) -> ()
+// CHECK-NEXT: acc.yield
+// CHECK-NEXT: }
+//
+// CHECK-NEXT: acc.private.recipe @privatization__ZTS13CopyConstruct : !cir.ptr<!rec_CopyConstruct> init {
+// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!rec_CopyConstruct> {{.*}}):
+// CHECK-NEXT: cir.alloca !rec_CopyConstruct, !cir.ptr<!rec_CopyConstruct>, ["openacc.private.init"]
+// CHECK-NEXT: acc.yield
+// CHECK-NEXT: }
+
+template<typename T, typename U, typename V, typename W>
+void dependent_version(const T &cc, const U &ndc, const V &dtor, const W &someInt) {
+ // CHECK: cir.func {{.*}}@_Z17dependent_versionI13CopyConstruct14NonDefaultCtor7HasDtoriEvRKT_RKT0_RKT1_RKT2_(%[[ARG0:.*]]: !cir.ptr<!rec_CopyConstruct> {{.*}}, %[[ARG1:.*]]: !cir.ptr<!rec_NonDefaultCtor> {{.*}}, %[[ARG2:.*]]: !cir.ptr<!rec_HasDtor> {{.*}}, %[[ARG3:.*]]: !cir.ptr<!s32i> {{.*}}) {
+ // CHECK-NEXT: %[[CC:.*]] = cir.alloca !cir.ptr<!rec_CopyConstruct>, !cir.ptr<!cir.ptr<!rec_CopyConstruct>>, ["cc", init, const]
+ // CHECK-NEXT: %[[NDC:.*]] = cir.alloca !cir.ptr<!rec_NonDefaultCtor>, !cir.ptr<!cir.ptr<!rec_NonDefaultCtor>>, ["ndc", init, const]
+ // CHECK-NEXT: %[[DTOR:.*]] = cir.alloca !cir.ptr<!rec_HasDtor>, !cir.ptr<!cir.ptr<!rec_HasDtor>>, ["dtor", init, const]
+ // CHECK-NEXT: %[[SOMEINT:.*]] = cir.alloca !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>, ["someInt", init, const]
+ // % 3 = cir.alloca !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>, ["someInt", init, const]
+
+#pragma acc parallel private(cc, ndc, dtor, someInt)
+ ;
+ // CHECK: %[[PRIV_LOAD:.*]] = cir.load %[[CC]] : !cir.ptr<!cir.ptr<!rec_CopyConstruct>>, !cir.ptr<!rec_CopyConstruct>
+ // CHECK-NEXT: %[[PRIVATE1:.*]] = acc.private varPtr(%[[PRIV_LOAD]] : !cir.ptr<!rec_CopyConstruct>) -> !cir.ptr<!rec_CopyConstruct> {name = "cc"}
+ // CHECK-NEXT: %[[PRIV_LOAD:.*]] = cir.load %[[NDC]] : !cir.ptr<!cir.ptr<!rec_NonDefaultCtor>>, !cir.ptr<!rec_NonDefaultCtor>
+ // CHECK-NEXT: %[[PRIVATE2:.*]] = acc.private varPtr(%[[PRIV_LOAD]] : !cir.ptr<!rec_NonDefaultCtor>) -> !cir.ptr<!rec_NonDefaultCtor> {name = "ndc"}
+ // CHECK-NEXT: %[[PRIV_LOAD:.*]] = cir.load %[[DTOR]] : !cir.ptr<!cir.ptr<!rec_HasDtor>>, !cir.ptr<!rec_HasDtor>
+ // CHECK-NEXT: %[[PRIVATE3:.*]] = acc.private varPtr(%[[PRIV_LOAD]] : !cir.ptr<!rec_HasDtor>) -> !cir.ptr<!rec_HasDtor> {name = "dtor"}
+ // CHECK-NEXT: %[[PRIV_LOAD:.*]] = cir.load %[[SOMEINT]] : !cir.ptr<!cir.ptr<!s32i>>, !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[PRIVATE4:.*]] = acc.private varPtr(%[[PRIV_LOAD]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "someInt"}
+
+ // CHECK-NEXT: acc.parallel private(@privatization__ZTS13CopyConstruct -> %[[PRIVATE1]] : !cir.ptr<!rec_CopyConstruct>,
+ // CHECK-SAME: @privatization__ZTS14NonDefaultCtor -> %[[PRIVATE2]] : !cir.ptr<!rec_NonDefaultCtor>,
+ // CHECK-SAME: @privatization__ZTS7HasDtor -> %[[PRIVATE3]] : !cir.ptr<!rec_HasDtor>,
+ // CHECK-SAME: @privatization__ZTSi -> %[[PRIVATE4]] : !cir.ptr<!s32i>) {
+ // CHECK-NEXT: acc.yield
+ // CHECK-NEXT: } loc
+}
+
+void use() {
+ CopyConstruct cc;
+ NonDefaultCtor ndc;
+ HasDtor dtor;
+ int i;
+ dependent_version(cc, ndc, dtor, i);
+}
diff --git a/clang/test/CIR/CodeGenOpenACC/compute-private-clause.cpp b/clang/test/CIR/CodeGenOpenACC/compute-private-clause.cpp
index a204f41ef7dbe..940bf53e7ece4 100644
--- a/clang/test/CIR/CodeGenOpenACC/compute-private-clause.cpp
+++ b/clang/test/CIR/CodeGenOpenACC/compute-private-clause.cpp
@@ -43,7 +43,24 @@ struct HasDtor {
//
// CHECK-NEXT: acc.private.recipe @privatization__ZTSA5_14NonDefaultCtor : !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>> init {
// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>> {{.*}}):
-// CHECK-NEXT: cir.alloca !cir.array<!rec_NonDefaultCtor x 5>, !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>>, ["openacc.private.init"]
+// CHECK-NEXT: %[[ALLOCA:.*]] = cir.alloca !cir.array<!rec_NonDefaultCtor x 5>, !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>>, ["openacc.private.init", init]
+// CHECK-NEXT: %[[LAST_IDX:.*]] = cir.const #cir.int<5> : !u64i
+// CHECK-NEXT: %[[ARRPTR:.*]] = cir.cast(array_to_ptrdecay, %[[ALLOCA]] : !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>>), !cir.ptr<!rec_NonDefaultCtor>
+// CHECK-NEXT: %[[LAST_ELEM:.*]] = cir.ptr_stride(%[[ARRPTR]] : !cir.ptr<!rec_NonDefaultCtor>, %[[LAST_IDX]] : !u64i), !cir.ptr<!rec_NonDefaultCtor>
+// CHECK-NEXT: %[[ITR:.*]] = cir.alloca !cir.ptr<!rec_NonDefaultCtor>, !cir.ptr<!cir.ptr<!rec_NonDefaultCtor>>, ["__array_idx"]
+// CHECK-NEXT: cir.store %[[ARRPTR]], %[[ITR]] : !cir.ptr<!rec_NonDefaultCtor>, !cir.ptr<!cir.ptr<!rec_NonDefaultCtor>>
+// CHECK-NEXT: cir.do {
+// CHECK-NEXT: %[[ELEM_LOAD:.*]] = cir.load %[[ITR]] : !cir.ptr<!cir.ptr<!rec_NonDefaultCtor>>, !cir.ptr<!rec_NonDefaultCtor>
+// CHECK-NEXT: cir.call @_ZN14NonDefaultCtorC1Ev(%[[ELEM_LOAD]]) : (!cir.ptr<!rec_NonDefaultCtor>) -> ()
+// CHECK-NEXT: %[[ONE_CONST:.*]] = cir.const #cir.int<1> : !u64i
+// CHECK-NEXT: %[[ELEM:.*]] = cir.ptr_stride(%[[ELEM_LOAD]] : !cir.ptr<!rec_NonDefaultCtor>, %[[ONE_CONST]] : !u64i), !cir.ptr<!rec_NonDefaultCtor>
+// CHECK-NEXT: cir.store %[[ELEM]], %[[ITR]] : !cir.ptr<!rec_NonDefaultCtor>, !cir.ptr<!cir.ptr<!rec_NonDefaultCtor>>
+// CHECK-NEXT: cir.yield
+// CHECK-NEXT: } while {
+// CHECK-NEXT: %[[ELEM_LOAD:.*]] = cir.load %[[ITR]] : !cir.ptr<!cir.ptr<!rec_NonDefaultCtor>>, !cir.ptr<!rec_NonDefaultCtor>
+// CHECK-NEXT: %[[CMP:.*]] = cir.cmp(ne, %[[ELEM_LOAD]], %[[LAST_ELEM]]) : !cir.ptr<!rec_NonDefaultCtor>, !cir.bool
+// CHECK-NEXT: cir.condition(%[[CMP]])
+// CHECK-NEXT: }
// CHECK-NEXT: acc.yield
// CHECK-NEXT: }
//
@@ -83,7 +100,8 @@ struct HasDtor {
//
// CHECK-NEXT: acc.private.recipe @privatization__ZTS14NonDefaultCtor : !cir.ptr<!rec_NonDefaultCtor> init {
// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!rec_NonDefaultCtor> {{.*}}):
-// CHECK-NEXT: cir.alloca !rec_NonDefaultCtor, !cir.ptr<!rec_NonDefaultCtor>, ["openacc.private.init"]
+// CHECK-NEXT: %[[ALLOCA:.*]] = cir.alloca !rec_NonDefaultCtor, !cir.ptr<!rec_NonDefaultCtor>, ["openacc.private.init", init]
+// CHECK-NEXT: cir.call @_ZN14NonDefaultCtorC1Ev(%[[ALLOCA]]) : (!cir.ptr<!rec_NonDefaultCtor>) -> ()
// CHECK-NEXT: acc.yield
// CHECK-NEXT: }
//
diff --git a/clang/test/CIR/CodeGenOpenACC/loop-private-clause.cpp b/clang/test/CIR/CodeGenOpenACC/loop-private-clause.cpp
index 384496bf87945..435a0c9157603 100644
--- a/clang/test/CIR/CodeGenOpenACC/loop-private-clause.cpp
+++ b/clang/test/CIR/CodeGenOpenACC/loop-private-clause.cpp
@@ -43,7 +43,24 @@ struct HasDtor {
//
// CHECK-NEXT: acc.private.recipe @privatization__ZTSA5_14NonDefaultCtor : !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>> init {
// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>> {{.*}}):
-// CHECK-NEXT: cir.alloca !cir.array<!rec_NonDefaultCtor x 5>, !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>>, ["openacc.private.init"]
+// CHECK-NEXT: %[[ALLOCA:.*]] = cir.alloca !cir.array<!rec_NonDefaultCtor x 5>, !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>>, ["openacc.private.init", init]
+// CHECK-NEXT: %[[LAST_IDX:.*]] = cir.const #cir.int<5> : !u64i
+// CHECK-NEXT: %[[ARRPTR:.*]] = cir.cast(array_to_ptrdecay, %[[ALLOCA]] : !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>>), !cir.ptr<!rec_NonDefaultCtor>
+// CHECK-NEXT: %[[LAST_ELEM:.*]] = cir.ptr_stride(%[[ARRPTR]] : !cir.ptr<!rec_NonDefaultCtor>, %[[LAST_IDX]] : !u64i), !cir.ptr<!rec_NonDefaultCtor>
+// CHECK-NEXT: %[[ITR:.*]] = cir.alloca !cir.ptr<!rec_NonDefaultCtor>, !cir.ptr<!cir.ptr<!rec_NonDefaultCtor>>, ["__array_idx"]
+// CHECK-NEXT: cir.store %[[ARRPTR]], %[[ITR]] : !cir.ptr<!rec_NonDefaultCtor>, !cir.ptr<!cir.ptr<!rec_NonDefaultCtor>>
+// CHECK-NEXT: cir.do {
+// CHECK-NEXT: %[[ELEM_LOAD:.*]] = cir.load %[[ITR]] : !cir.ptr<!cir.ptr<!rec_NonDefaultCtor>>, !cir.ptr<!rec_NonDefaultCtor>
+// CHECK-NEXT: cir.call @_ZN14NonDefaultCtorC1Ev(%[[ELEM_LOAD]]) : (!cir.ptr<!rec_NonDefaultCtor>) -> ()
+// CHECK-NEXT: %[[ONE_CONST:.*]] = cir.const #cir.int<1> : !u64i
+// CHECK-NEXT: %[[ELEM:.*]] = cir.ptr_stride(%[[ELEM_LOAD]] : !cir.ptr<!rec_NonDefaultCtor>, %[[ONE_CONST]] : !u64i), !cir.ptr<!rec_NonDefaultCtor>
+// CHECK-NEXT: cir.store %[[ELEM]], %[[ITR]] : !cir.ptr<!rec_NonDefaultCtor>, !cir.ptr<!cir.ptr<!rec_NonDefaultCtor>>
+// CHECK-NEXT: cir.yield
+// CHECK-NEXT: } while {
+// CHECK-NEXT: %[[ELEM_LOAD:.*]] = cir.load %[[ITR]] : !cir.ptr<!cir.ptr<!rec_NonDefaultCtor>>, !cir.ptr<!rec_NonDefaultCtor>
+// CHECK-NEXT: %[[CMP:.*]] = cir.cmp(ne, %[[ELEM_LOAD]], %[[LAST_ELEM]]) : !cir.ptr<!rec_NonDefaultCtor>, !cir.bool
+// CHECK-NEXT: cir.condition(%[[CMP]])
+// CHECK-NEXT: }
// CHECK-NEXT: acc.yield
// CHECK-NEXT: }
//
@@ -83,7 +100,8 @@ struct HasDtor {
//
// CHECK-NEXT: acc.private.recipe @privatization__ZTS14NonDefaultCtor : !cir.ptr<!rec_NonDefaultCtor> init {
// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!rec_NonDefaultCtor> {{.*}}):
-// CHECK-NEXT: cir.alloca !rec_NonDefaultCtor, !cir.ptr<!rec_NonDefaultCtor>, ["openacc.private.init"]
+// CHECK-NEXT: %[[ALLOCA:.*]] = cir.alloca !rec_NonDefaultCtor, !cir.ptr<!rec_NonDefaultCtor>, ["openacc.private.init", init]
+// CHECK-NEXT: cir.call @_ZN14NonDefaultCtorC1Ev(%[[ALLOCA]]) : (!cir.ptr<!rec_NonDefaultCtor>) -> ()
// CHECK-NEXT: acc.yield
// CHECK-NEXT: }
//
diff --git a/clang/tools/libclang/CIndex.cpp b/clang/tools/libclang/CIndex.cpp
index 9412d9735ef82..9c2724f17cf8b 100644
--- a/clang/tools/libclang/CIndex.cpp
+++ b/clang/tools/libclang/CIndex.cpp
@@ -2880,6 +2880,8 @@ void OpenACCClauseEnqueue::VisitTileClause(const OpenACCTileClause &C) {
void OpenACCClauseEnqueue::VisitPrivateClause(const OpenACCPrivateClause &C) {
VisitVarList(C);
+ for (VarDecl *V : C.getInitRecipes())
+ Visitor.AddDecl(V);
}
void OpenACCClauseEnqueue::VisitHostClause(const OpenACCHostClause &C) {
More information about the cfe-commits
mailing list