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) { _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
