Author: erichkeane Date: 2024-12-05T06:35:36-08:00 New Revision: 3a4b9f38915625c68c78b62de48a3de8b97c5043
URL: https://github.com/llvm/llvm-project/commit/3a4b9f38915625c68c78b62de48a3de8b97c5043 DIFF: https://github.com/llvm/llvm-project/commit/3a4b9f38915625c68c78b62de48a3de8b97c5043.diff LOG: [OpenACC] Implement 'gang' clause for Combined Constructs This one is a bit complicated, as it has some interesting interactions, as 'gang' Sema is required to look at its containing compute construct. Except in the case of a combined construct, they are the same. This resulted in a large refactor of the checking code for CheckGangExpr, plus some additional work on the diagnostics for its interaction with 'num_gangs' and 'vector'/'worker'. Added: clang/test/SemaOpenACC/combined-construct-gang-ast.cpp clang/test/SemaOpenACC/combined-construct-gang-clause.cpp Modified: clang/include/clang/AST/OpenACCClause.h clang/include/clang/Basic/DiagnosticSemaKinds.td clang/include/clang/Sema/SemaOpenACC.h clang/lib/Sema/SemaOpenACC.cpp clang/lib/Sema/TreeTransform.h clang/test/AST/ast-print-openacc-combined-construct.cpp clang/test/SemaOpenACC/combined-construct-auto_seq_independent-clauses.c clang/test/SemaOpenACC/combined-construct-device_type-clause.c clang/test/SemaOpenACC/loop-construct-gang-clause.cpp Removed: ################################################################################ diff --git a/clang/include/clang/AST/OpenACCClause.h b/clang/include/clang/AST/OpenACCClause.h index 5ad4c336b6c531..2588c3f645c02b 100644 --- a/clang/include/clang/AST/OpenACCClause.h +++ b/clang/include/clang/AST/OpenACCClause.h @@ -483,6 +483,14 @@ class OpenACCGangClause final return {getGangKind(I), getExprs()[I]}; } + bool hasExprOfKind(OpenACCGangKind GK) const { + for (unsigned I = 0; I < getNumExprs(); ++I) { + if (getGangKind(I) == GK) + return true; + } + return false; + } + static OpenACCGangClause * Create(const ASTContext &Ctx, SourceLocation BeginLoc, SourceLocation LParenLoc, ArrayRef<OpenACCGangKind> GangKinds, diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 2137cb713164ad..447358f0a5f382 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -12755,21 +12755,22 @@ def err_acc_gang_multiple_elt : Error<"OpenACC 'gang' clause may have at most one %select{unnamed or " "'num'|'dim'|'static'}0 argument">; def err_acc_int_arg_invalid - : Error<"'%1' argument on '%0' clause is not permitted on a%select{n " - "orphaned|||}2 'loop' construct %select{|associated with a " - "'parallel' compute construct|associated with a 'kernels' compute " - "construct|associated with a 'serial' compute construct}2">; + : Error<"'%0' argument on '%1' clause is not permitted on a%select{|n " + "orphaned}2 '%3' construct%select{| associated with a '%5' compute " + "construct}4">; def err_acc_gang_dim_value : Error<"argument to 'gang' clause dimension must be %select{a constant " "expression|1, 2, or 3: evaluated to %1}0">; def err_acc_num_arg_conflict - : Error<"'num' argument to '%0' clause not allowed on a 'loop' construct " - "associated with a 'kernels' construct that has a " - "'%select{num_gangs|num_workers|vector_length}1' " - "clause">; + : Error<"'num' argument to '%0' clause not allowed on a '%1' " + "construct%select{| associated with a '%3' construct}2 that has a " + "'%4' clause">; +def err_acc_num_arg_conflict_reverse + : Error<"'num_gangs' clause not allowed on a 'kernels loop' construct that " + "has a 'gang' clause with a 'num' argument">; def err_acc_clause_in_clause_region : Error<"loop with a '%0' clause may not exist in the region of a '%1' " - "clause%select{| on a 'kernels' compute construct}2">; + "clause%select{| on a '%3' construct}2">; def err_acc_gang_reduction_conflict : Error<"%select{OpenACC 'gang' clause with a 'dim' value greater than " "1|OpenACC 'reduction' clause}0 cannot " diff --git a/clang/include/clang/Sema/SemaOpenACC.h b/clang/include/clang/Sema/SemaOpenACC.h index d720cf3c74d87e..a4132534686a81 100644 --- a/clang/include/clang/Sema/SemaOpenACC.h +++ b/clang/include/clang/Sema/SemaOpenACC.h @@ -164,9 +164,14 @@ class SemaOpenACC : public SemaBase { } /// If there is a current 'active' loop construct with a 'gang' clause on a - /// 'kernel' construct, this will have the source location for it. This - /// permits us to implement the restriction of no further 'gang' clauses. - SourceLocation LoopGangClauseOnKernelLoc; + /// 'kernel' construct, this will have the source location for it, and the + /// 'kernel kind'. This permits us to implement the restriction of no further + /// 'gang' clauses. + struct LoopGangOnKernelTy { + SourceLocation Loc; + OpenACCDirectiveKind DirKind = OpenACCDirectiveKind::Invalid; + } LoopGangClauseOnKernel; + /// If there is a current 'active' loop construct with a 'worker' clause on it /// (on any sort of construct), this has the source location for it. This /// permits us to implement the restriction of no further 'gang' or 'worker' @@ -705,7 +710,9 @@ class SemaOpenACC : public SemaBase { ExprResult CheckTileSizeExpr(Expr *SizeExpr); // Check a single expression on a gang clause. - ExprResult CheckGangExpr(OpenACCGangKind GK, Expr *E); + ExprResult CheckGangExpr(ArrayRef<const OpenACCClause *> ExistingClauses, + OpenACCDirectiveKind DK, OpenACCGangKind GK, + Expr *E); // Does the checking for a 'gang' clause that needs to be done in dependent // and not dependent cases. @@ -771,7 +778,7 @@ class SemaOpenACC : public SemaBase { SemaOpenACC &SemaRef; ComputeConstructInfo OldActiveComputeConstructInfo; OpenACCDirectiveKind DirKind; - SourceLocation OldLoopGangClauseOnKernelLoc; + LoopGangOnKernelTy OldLoopGangClauseOnKernel; SourceLocation OldLoopWorkerClauseLoc; SourceLocation OldLoopVectorClauseLoc; LoopWithoutSeqCheckingInfo OldLoopWithoutSeqInfo; diff --git a/clang/lib/Sema/SemaOpenACC.cpp b/clang/lib/Sema/SemaOpenACC.cpp index 30491d6f3a3bd6..16348d0b8837ef 100644 --- a/clang/lib/Sema/SemaOpenACC.cpp +++ b/clang/lib/Sema/SemaOpenACC.cpp @@ -737,6 +737,24 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitNumGangsClause( return nullptr; } } + + // OpenACC 3.3 Section 2.9.2: + // An argument with no keyword or with the 'num' wkeyword is allowed only when + // the 'num_gangs' does not appear on the 'kernel' construct. + if (Clause.getDirectiveKind() == OpenACCDirectiveKind::KernelsLoop) { + auto GangClauses = llvm::make_filter_range( + ExistingClauses, llvm::IsaPred<OpenACCGangClause>); + + for (auto *GC : GangClauses) { + if (cast<OpenACCGangClause>(GC)->hasExprOfKind(OpenACCGangKind::Num)) { + SemaRef.Diag(Clause.getBeginLoc(), + diag::err_acc_num_arg_conflict_reverse); + SemaRef.Diag(GC->getBeginLoc(), diag::note_acc_previous_clause_here); + return nullptr; + } + } + } + return OpenACCNumGangsClause::Create( Ctx, Clause.getBeginLoc(), Clause.getLParenLoc(), Clause.getIntExprs(), Clause.getEndLoc()); @@ -1033,6 +1051,136 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitIndependentClause( Clause.getEndLoc()); } +ExprResult CheckGangStaticExpr(SemaOpenACC &S, Expr *E) { + if (isa<OpenACCAsteriskSizeExpr>(E)) + return E; + return S.ActOnIntExpr(OpenACCDirectiveKind::Invalid, OpenACCClauseKind::Gang, + E->getBeginLoc(), E); +} + +bool IsOrphanLoop(OpenACCDirectiveKind DK, OpenACCDirectiveKind AssocKind) { + return DK == OpenACCDirectiveKind::Loop && + AssocKind == OpenACCDirectiveKind::Invalid; +} + +bool HasAssocKind(OpenACCDirectiveKind DK, OpenACCDirectiveKind AssocKind) { + return DK == OpenACCDirectiveKind::Loop && + AssocKind != OpenACCDirectiveKind::Invalid; +} + +ExprResult DiagIntArgInvalid(SemaOpenACC &S, Expr *E, OpenACCGangKind GK, + OpenACCClauseKind CK, OpenACCDirectiveKind DK, + OpenACCDirectiveKind AssocKind) { + S.Diag(E->getBeginLoc(), diag::err_acc_int_arg_invalid) + << GK << CK << IsOrphanLoop(DK, AssocKind) << DK + << HasAssocKind(DK, AssocKind) << AssocKind; + return ExprError(); +} + +ExprResult CheckGangParallelExpr(SemaOpenACC &S, OpenACCDirectiveKind DK, + OpenACCDirectiveKind AssocKind, + OpenACCGangKind GK, Expr *E) { + switch (GK) { + case OpenACCGangKind::Static: + return CheckGangStaticExpr(S, E); + case OpenACCGangKind::Num: + // OpenACC 3.3 2.9.2: When the parent compute construct is a parallel + // construct, or an orphaned loop construct, the gang clause behaves as + // follows. ... The num argument is not allowed. + return DiagIntArgInvalid(S, E, GK, OpenACCClauseKind::Gang, DK, AssocKind); + case OpenACCGangKind::Dim: { + // OpenACC 3.3 2.9.2: When the parent compute construct is a parallel + // construct, or an orphaned loop construct, the gang clause behaves as + // follows. ... The dim argument must be a constant positive integer value + // 1, 2, or 3. + if (!E) + return ExprError(); + ExprResult Res = + S.ActOnIntExpr(OpenACCDirectiveKind::Invalid, OpenACCClauseKind::Gang, + E->getBeginLoc(), E); + + if (!Res.isUsable()) + return Res; + + if (Res.get()->isInstantiationDependent()) + return Res; + + std::optional<llvm::APSInt> ICE = + Res.get()->getIntegerConstantExpr(S.getASTContext()); + + if (!ICE || *ICE <= 0 || ICE > 3) { + S.Diag(Res.get()->getBeginLoc(), diag::err_acc_gang_dim_value) + << ICE.has_value() << ICE.value_or(llvm::APSInt{}).getExtValue(); + return ExprError(); + } + + return ExprResult{ + ConstantExpr::Create(S.getASTContext(), Res.get(), APValue{*ICE})}; + } + } + llvm_unreachable("Unknown gang kind in gang parallel check"); +} + +ExprResult CheckGangKernelsExpr(SemaOpenACC &S, + ArrayRef<const OpenACCClause *> ExistingClauses, + OpenACCDirectiveKind DK, + OpenACCDirectiveKind AssocKind, + OpenACCGangKind GK, Expr *E) { + switch (GK) { + // OpenACC 3.3 2.9.2: When the parent compute construct is a kernels + // construct, the gang clause behaves as follows. ... The dim argument is + // not allowed. + case OpenACCGangKind::Dim: + return DiagIntArgInvalid(S, E, GK, OpenACCClauseKind::Gang, DK, AssocKind); + case OpenACCGangKind::Num: { + // OpenACC 3.3 2.9.2: When the parent compute construct is a kernels + // construct, the gang clause behaves as follows. ... An argument with no + // keyword or with num keyword is only allowed when num_gangs does not + // appear on the kernels construct. ... The region of a loop with the gang + // clause may not contain another loop with a gang clause unless within a + // nested compute region. + + // If this is a 'combined' construct, search the list of existing clauses. + // Else we need to search the containing 'kernel'. + auto Collection = isOpenACCCombinedDirectiveKind(DK) + ? ExistingClauses + : S.getActiveComputeConstructInfo().Clauses; + + const auto *Itr = + llvm::find_if(Collection, llvm::IsaPred<OpenACCNumGangsClause>); + + if (Itr != Collection.end()) { + S.Diag(E->getBeginLoc(), diag::err_acc_num_arg_conflict) + << OpenACCClauseKind::Gang << DK << HasAssocKind(DK, AssocKind) + << AssocKind << OpenACCClauseKind::NumGangs; + + S.Diag((*Itr)->getBeginLoc(), diag::note_acc_previous_clause_here); + return ExprError(); + } + return ExprResult{E}; + } + case OpenACCGangKind::Static: + return CheckGangStaticExpr(S, E); + return ExprError(); + } + llvm_unreachable("Unknown gang kind in gang kernels check"); +} + +ExprResult CheckGangSerialExpr(SemaOpenACC &S, OpenACCDirectiveKind DK, + OpenACCDirectiveKind AssocKind, + OpenACCGangKind GK, Expr *E) { + switch (GK) { + // 'dim' and 'num' don't really make sense on serial, and GCC rejects them + // too, so we disallow them too. + case OpenACCGangKind::Dim: + case OpenACCGangKind::Num: + return DiagIntArgInvalid(S, E, GK, OpenACCClauseKind::Gang, DK, AssocKind); + case OpenACCGangKind::Static: + return CheckGangStaticExpr(S, E); + } + llvm_unreachable("Unknown gang kind in gang serial check"); +} + OpenACCClause *SemaOpenACCClauseVisitor::VisitVectorClause( SemaOpenACC::OpenACCParsedClause &Clause) { if (DiagIfSeqClause(Clause)) @@ -1054,8 +1202,9 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitVectorClause( case OpenACCDirectiveKind::Serial: // GCC disallows this, and there is no real good reason for us to permit // it, so disallow until we come up with a use case that makes sense. - SemaRef.Diag(IntExpr->getBeginLoc(), diag::err_acc_int_arg_invalid) - << OpenACCClauseKind::Vector << "num" << /*serial=*/3; + DiagIntArgInvalid(SemaRef, IntExpr, OpenACCGangKind::Num, + OpenACCClauseKind::Vector, Clause.getDirectiveKind(), + SemaRef.getActiveComputeConstructInfo().Kind); IntExpr = nullptr; break; case OpenACCDirectiveKind::Kernels: { @@ -1064,7 +1213,11 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitVectorClause( llvm::IsaPred<OpenACCVectorLengthClause>); if (Itr != SemaRef.getActiveComputeConstructInfo().Clauses.end()) { SemaRef.Diag(IntExpr->getBeginLoc(), diag::err_acc_num_arg_conflict) - << OpenACCClauseKind::Vector << /*vector_length=*/2; + << OpenACCClauseKind::Vector << Clause.getDirectiveKind() + << HasAssocKind(Clause.getDirectiveKind(), + SemaRef.getActiveComputeConstructInfo().Kind) + << SemaRef.getActiveComputeConstructInfo().Kind + << OpenACCClauseKind::VectorLength; SemaRef.Diag((*Itr)->getBeginLoc(), diag::note_acc_previous_clause_here); @@ -1113,18 +1266,11 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitWorkerClause( if (IntExpr) { switch (SemaRef.getActiveComputeConstructInfo().Kind) { case OpenACCDirectiveKind::Invalid: - SemaRef.Diag(IntExpr->getBeginLoc(), diag::err_acc_int_arg_invalid) - << OpenACCClauseKind::Worker << "num" << /*orphan=*/0; - IntExpr = nullptr; - break; case OpenACCDirectiveKind::Parallel: - SemaRef.Diag(IntExpr->getBeginLoc(), diag::err_acc_int_arg_invalid) - << OpenACCClauseKind::Worker << "num" << /*parallel=*/1; - IntExpr = nullptr; - break; case OpenACCDirectiveKind::Serial: - SemaRef.Diag(IntExpr->getBeginLoc(), diag::err_acc_int_arg_invalid) - << OpenACCClauseKind::Worker << "num" << /*serial=*/3; + DiagIntArgInvalid(SemaRef, IntExpr, OpenACCGangKind::Num, + OpenACCClauseKind::Worker, Clause.getDirectiveKind(), + SemaRef.getActiveComputeConstructInfo().Kind); IntExpr = nullptr; break; case OpenACCDirectiveKind::Kernels: { @@ -1133,7 +1279,11 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitWorkerClause( llvm::IsaPred<OpenACCNumWorkersClause>); if (Itr != SemaRef.getActiveComputeConstructInfo().Clauses.end()) { SemaRef.Diag(IntExpr->getBeginLoc(), diag::err_acc_num_arg_conflict) - << OpenACCClauseKind::Worker << /*num_workers=*/1; + << OpenACCClauseKind::Worker << Clause.getDirectiveKind() + << HasAssocKind(Clause.getDirectiveKind(), + SemaRef.getActiveComputeConstructInfo().Kind) + << SemaRef.getActiveComputeConstructInfo().Kind + << OpenACCClauseKind::NumWorkers; SemaRef.Diag((*Itr)->getBeginLoc(), diag::note_acc_previous_clause_here); @@ -1187,12 +1337,15 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitGangClause( // Restrictions only properly implemented on 'loop' constructs, and it is // the only construct that can do anything with this, so skip/treat as // unimplemented for the combined constructs. - if (Clause.getDirectiveKind() != OpenACCDirectiveKind::Loop) + if (Clause.getDirectiveKind() != OpenACCDirectiveKind::Loop && + !isOpenACCCombinedDirectiveKind(Clause.getDirectiveKind())) return isNotImplemented(); // OpenACC 3.3 Section 2.9.11: A reduction clause may not appear on a loop // directive that has a gang clause and is within a compute construct that has // a num_gangs clause with more than one explicit argument. + // TODO OpenACC: When we implement reduction on combined constructs, we need + // to do this too. if (Clause.getDirectiveKind() == OpenACCDirectiveKind::Loop && SemaRef.getActiveComputeConstructInfo().Kind != OpenACCDirectiveKind::Invalid) { @@ -1229,31 +1382,13 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitGangClause( for (unsigned I = 0; I < Clause.getIntExprs().size(); ++I) { OpenACCGangKind GK = Clause.getGangKinds()[I]; - ExprResult ER = SemaRef.CheckGangExpr(GK, Clause.getIntExprs()[I]); + ExprResult ER = + SemaRef.CheckGangExpr(ExistingClauses, Clause.getDirectiveKind(), GK, + Clause.getIntExprs()[I]); if (!ER.isUsable()) continue; - // OpenACC 3.3 2.9.2: When the parent compute construct is a kernels - // construct, the gang clause behaves as follows. ... An argument with no - // keyword or with num keyword is only allowed when num_gangs does not - // appear on the kernels construct. - if (SemaRef.getActiveComputeConstructInfo().Kind == - OpenACCDirectiveKind::Kernels && - GK == OpenACCGangKind::Num) { - const auto *Itr = - llvm::find_if(SemaRef.getActiveComputeConstructInfo().Clauses, - llvm::IsaPred<OpenACCNumGangsClause>); - - if (Itr != SemaRef.getActiveComputeConstructInfo().Clauses.end()) { - SemaRef.Diag(ER.get()->getBeginLoc(), diag::err_acc_num_arg_conflict) - << OpenACCClauseKind::Gang << /*num_gangs=*/0; - SemaRef.Diag((*Itr)->getBeginLoc(), - diag::note_acc_previous_clause_here); - continue; - } - } - // OpenACC 3.3 2.9: 'gang-arg-list' may have at most one num, one dim, and // one static argument. if (ExistingElemLoc[static_cast<unsigned>(GK)].isValid()) { @@ -1269,47 +1404,50 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitGangClause( IntExprs.push_back(ER.get()); } - // OpenACC 3.3 2.9.2: When the parent compute construct is a kernels - // construct, the gang clause behaves as follows. ... The region of a loop - // with a gang clause may not contain another loop with a gang clause unless - // within a nested compute region. - if (SemaRef.LoopGangClauseOnKernelLoc.isValid()) { - // This handles the 'inner loop' diagnostic, but we cannot set that we're on - // one of these until we get to the end of the construct. - SemaRef.Diag(Clause.getBeginLoc(), diag::err_acc_clause_in_clause_region) - << OpenACCClauseKind::Gang << OpenACCClauseKind::Gang - << /*kernels construct info*/ 1; - SemaRef.Diag(SemaRef.LoopGangClauseOnKernelLoc, - diag::note_acc_previous_clause_here); - return nullptr; - } + if (!isOpenACCCombinedDirectiveKind(Clause.getDirectiveKind())) { + // OpenACC 3.3 2.9.2: When the parent compute construct is a kernels + // construct, the gang clause behaves as follows. ... The region of a loop + // with a gang clause may not contain another loop with a gang clause unless + // within a nested compute region. + if (SemaRef.LoopGangClauseOnKernel.Loc.isValid()) { + // This handles the 'inner loop' diagnostic, but we cannot set that we're + // on one of these until we get to the end of the construct. + SemaRef.Diag(Clause.getBeginLoc(), diag::err_acc_clause_in_clause_region) + << OpenACCClauseKind::Gang << OpenACCClauseKind::Gang + << /*kernels construct info*/ 1 + << SemaRef.LoopGangClauseOnKernel.DirKind; + SemaRef.Diag(SemaRef.LoopGangClauseOnKernel.Loc, + diag::note_acc_previous_clause_here); + return nullptr; + } - // OpenACC 3.3 2.9.3: The region of a loop with a 'worker' clause may not - // contain a loop with a gang or worker clause unless within a nested compute - // region. - if (SemaRef.LoopWorkerClauseLoc.isValid()) { - // This handles the 'inner loop' diagnostic, but we cannot set that we're on - // one of these until we get to the end of the construct. - SemaRef.Diag(Clause.getBeginLoc(), diag::err_acc_clause_in_clause_region) - << OpenACCClauseKind::Gang << OpenACCClauseKind::Worker - << /*kernels construct info*/ 1; - SemaRef.Diag(SemaRef.LoopWorkerClauseLoc, - diag::note_acc_previous_clause_here); - return nullptr; - } + // OpenACC 3.3 2.9.3: The region of a loop with a 'worker' clause may not + // contain a loop with a gang or worker clause unless within a nested + // compute region. + if (SemaRef.LoopWorkerClauseLoc.isValid()) { + // This handles the 'inner loop' diagnostic, but we cannot set that we're + // on one of these until we get to the end of the construct. + SemaRef.Diag(Clause.getBeginLoc(), diag::err_acc_clause_in_clause_region) + << OpenACCClauseKind::Gang << OpenACCClauseKind::Worker + << /*!kernels construct info*/ 0; + SemaRef.Diag(SemaRef.LoopWorkerClauseLoc, + diag::note_acc_previous_clause_here); + return nullptr; + } - // OpenACC 3.3 2.9.4: The region of a loop with a 'vector' clause may not - // contain a loop with a gang, worker, or vector clause unless within a nested - // compute region. - if (SemaRef.LoopVectorClauseLoc.isValid()) { - // This handles the 'inner loop' diagnostic, but we cannot set that we're on - // one of these until we get to the end of the construct. - SemaRef.Diag(Clause.getBeginLoc(), diag::err_acc_clause_in_clause_region) - << OpenACCClauseKind::Gang << OpenACCClauseKind::Vector - << /*kernels construct info*/ 1; - SemaRef.Diag(SemaRef.LoopVectorClauseLoc, - diag::note_acc_previous_clause_here); - return nullptr; + // OpenACC 3.3 2.9.4: The region of a loop with a 'vector' clause may not + // contain a loop with a gang, worker, or vector clause unless within a + // nested compute region. + if (SemaRef.LoopVectorClauseLoc.isValid()) { + // This handles the 'inner loop' diagnostic, but we cannot set that we're + // on one of these until we get to the end of the construct. + SemaRef.Diag(Clause.getBeginLoc(), diag::err_acc_clause_in_clause_region) + << OpenACCClauseKind::Gang << OpenACCClauseKind::Vector + << /*!kernels construct info*/ 0; + SemaRef.Diag(SemaRef.LoopVectorClauseLoc, + diag::note_acc_previous_clause_here); + return nullptr; + } } return SemaRef.CheckGangClause(ExistingClauses, Clause.getBeginLoc(), @@ -1489,7 +1627,7 @@ SemaOpenACC::AssociatedStmtRAII::AssociatedStmtRAII( ArrayRef<const OpenACCClause *> UnInstClauses, ArrayRef<OpenACCClause *> Clauses) : SemaRef(S), OldActiveComputeConstructInfo(S.ActiveComputeConstructInfo), - DirKind(DK), OldLoopGangClauseOnKernelLoc(S.LoopGangClauseOnKernelLoc), + DirKind(DK), OldLoopGangClauseOnKernel(S.LoopGangClauseOnKernel), OldLoopWorkerClauseLoc(S.LoopWorkerClauseLoc), OldLoopVectorClauseLoc(S.LoopVectorClauseLoc), OldLoopWithoutSeqInfo(S.LoopWithoutSeqInfo), @@ -1510,7 +1648,7 @@ SemaOpenACC::AssociatedStmtRAII::AssociatedStmtRAII( // within a nested compute region. // // Implement the 'unless within a nested compute region' part. - SemaRef.LoopGangClauseOnKernelLoc = {}; + SemaRef.LoopGangClauseOnKernel = {}; SemaRef.LoopWorkerClauseLoc = {}; SemaRef.LoopVectorClauseLoc = {}; SemaRef.LoopWithoutSeqInfo = {}; @@ -1524,8 +1662,7 @@ SemaOpenACC::AssociatedStmtRAII::AssociatedStmtRAII( SetCollapseInfoBeforeAssociatedStmt(UnInstClauses, Clauses); SetTileInfoBeforeAssociatedStmt(UnInstClauses, Clauses); - // TODO: OpenACC: We need to set these 3, CollapseInfo, and TileInfo - SemaRef.LoopGangClauseOnKernelLoc = {}; + SemaRef.LoopGangClauseOnKernel = {}; SemaRef.LoopWorkerClauseLoc = {}; SemaRef.LoopVectorClauseLoc = {}; @@ -1548,7 +1685,7 @@ SemaOpenACC::AssociatedStmtRAII::AssociatedStmtRAII( // This handles the 'outer loop' part of this. auto *Itr = llvm::find_if(Clauses, llvm::IsaPred<OpenACCGangClause>); if (Itr != Clauses.end()) - SemaRef.LoopGangClauseOnKernelLoc = (*Itr)->getBeginLoc(); + SemaRef.LoopGangClauseOnKernel = {(*Itr)->getBeginLoc(), DirKind}; } if (UnInstClauses.empty()) { @@ -1586,7 +1723,8 @@ SemaOpenACC::AssociatedStmtRAII::AssociatedStmtRAII( // This handles the 'outer loop' part of this. auto *Itr = llvm::find_if(Clauses, llvm::IsaPred<OpenACCGangClause>); if (Itr != Clauses.end()) - SemaRef.LoopGangClauseOnKernelLoc = (*Itr)->getBeginLoc(); + SemaRef.LoopGangClauseOnKernel = {(*Itr)->getBeginLoc(), + OpenACCDirectiveKind::Kernels}; } if (UnInstClauses.empty()) { @@ -1671,7 +1809,7 @@ void SemaOpenACC::AssociatedStmtRAII::SetTileInfoBeforeAssociatedStmt( SemaOpenACC::AssociatedStmtRAII::~AssociatedStmtRAII() { SemaRef.ActiveComputeConstructInfo = OldActiveComputeConstructInfo; - SemaRef.LoopGangClauseOnKernelLoc = OldLoopGangClauseOnKernelLoc; + SemaRef.LoopGangClauseOnKernel = OldLoopGangClauseOnKernel; SemaRef.LoopWorkerClauseLoc = OldLoopWorkerClauseLoc; SemaRef.LoopVectorClauseLoc = OldLoopVectorClauseLoc; SemaRef.LoopWithoutSeqInfo = OldLoopWithoutSeqInfo; @@ -2328,106 +2466,47 @@ ExprResult SemaOpenACC::CheckCollapseLoopCount(Expr *LoopCount) { ConstantExpr::Create(getASTContext(), LoopCount, APValue{*ICE})}; } -namespace { -ExprResult CheckGangStaticExpr(SemaOpenACC &S, Expr *E) { - if (isa<OpenACCAsteriskSizeExpr>(E)) - return E; - return S.ActOnIntExpr(OpenACCDirectiveKind::Invalid, OpenACCClauseKind::Gang, - E->getBeginLoc(), E); -} -} // namespace - -ExprResult SemaOpenACC::CheckGangExpr(OpenACCGangKind GK, Expr *E) { - // Gang Expr legality depends on the associated compute construct. - switch (ActiveComputeConstructInfo.Kind) { - case OpenACCDirectiveKind::Invalid: - case OpenACCDirectiveKind::Parallel: { - switch (GK) { - // OpenACC 3.3 2.9.2: When the parent compute construct is a parallel - // construct, or an orphaned loop construct, the gang clause behaves as - // follows. ... The dim argument must be a constant positive integer value - // 1, 2, or 3. - case OpenACCGangKind::Dim: { - if (!E) - return ExprError(); - ExprResult Res = - ActOnIntExpr(OpenACCDirectiveKind::Invalid, OpenACCClauseKind::Gang, - E->getBeginLoc(), E); - - if (!Res.isUsable()) - return Res; - - if (Res.get()->isInstantiationDependent()) - return Res; - - std::optional<llvm::APSInt> ICE = - Res.get()->getIntegerConstantExpr(getASTContext()); - - if (!ICE || *ICE <= 0 || ICE > 3) { - Diag(Res.get()->getBeginLoc(), diag::err_acc_gang_dim_value) - << ICE.has_value() << ICE.value_or(llvm::APSInt{}).getExtValue(); - return ExprError(); - } - - return ExprResult{ - ConstantExpr::Create(getASTContext(), Res.get(), APValue{*ICE})}; - } - // OpenACC 3.3 2.9.2: When the parent compute construct is a parallel - // construct, or an orphaned loop construct, the gang clause behaves as - // follows. ... The num argument is not allowed. - case OpenACCGangKind::Num: - Diag(E->getBeginLoc(), diag::err_acc_int_arg_invalid) - << OpenACCClauseKind::Gang << GK - << (/*orphan/parallel=*/ActiveComputeConstructInfo.Kind == - OpenACCDirectiveKind::Parallel - ? 1 - : 0); - return ExprError(); - case OpenACCGangKind::Static: - return CheckGangStaticExpr(*this, E); - } - } break; - case OpenACCDirectiveKind::Kernels: { - switch (GK) { - // OpenACC 3.3 2.9.2: When the parent compute construct is a kernels - // construct, the gang clause behaves as follows. ... The dim argument is - // not allowed. - case OpenACCGangKind::Dim: - Diag(E->getBeginLoc(), diag::err_acc_int_arg_invalid) - << OpenACCClauseKind::Gang << GK << /*kernels=*/2; - return ExprError(); - // OpenACC 3.3 2.9.2: When the parent compute construct is a kernels - // construct, the gang clause behaves as follows. ... An argument with no - // keyword or with num keyword is only allowed when num_gangs does not - // appear on the kernels construct. ... The region of a loop with the gang - // clause may not contain another loop with a gang clause unless within a - // nested compute region. - case OpenACCGangKind::Num: - // This isn't allowed if there is a 'num_gangs' on the kernel construct, - // and makes loop-with-gang-clause ill-formed inside of this 'loop', but - // nothing can be enforced here. - return ExprResult{E}; - case OpenACCGangKind::Static: - return CheckGangStaticExpr(*this, E); - } - } break; - case OpenACCDirectiveKind::Serial: { - switch (GK) { - // 'dim' and 'num' don't really make sense on serial, and GCC rejects them - // too, so we disallow them too. - case OpenACCGangKind::Dim: - case OpenACCGangKind::Num: - Diag(E->getBeginLoc(), diag::err_acc_int_arg_invalid) - << OpenACCClauseKind::Gang << GK << /*Kernels=*/3; - return ExprError(); - case OpenACCGangKind::Static: - return CheckGangStaticExpr(*this, E); +ExprResult +SemaOpenACC::CheckGangExpr(ArrayRef<const OpenACCClause *> ExistingClauses, + OpenACCDirectiveKind DK, OpenACCGangKind GK, + Expr *E) { + // There are two cases for the enforcement here: the 'current' directive is a + // 'loop', where we need to check the active compute construct kind, or the + // current directive is a 'combined' construct, where we have to check the + // current one. + switch (DK) { + case OpenACCDirectiveKind::ParallelLoop: + return CheckGangParallelExpr(*this, DK, ActiveComputeConstructInfo.Kind, GK, + E); + case OpenACCDirectiveKind::SerialLoop: + return CheckGangSerialExpr(*this, DK, ActiveComputeConstructInfo.Kind, GK, + E); + case OpenACCDirectiveKind::KernelsLoop: + return CheckGangKernelsExpr(*this, ExistingClauses, DK, + ActiveComputeConstructInfo.Kind, GK, E); + case OpenACCDirectiveKind::Loop: + switch (ActiveComputeConstructInfo.Kind) { + case OpenACCDirectiveKind::Invalid: + case OpenACCDirectiveKind::Parallel: + case OpenACCDirectiveKind::ParallelLoop: + return CheckGangParallelExpr(*this, DK, ActiveComputeConstructInfo.Kind, + GK, E); + case OpenACCDirectiveKind::SerialLoop: + case OpenACCDirectiveKind::Serial: + return CheckGangSerialExpr(*this, DK, ActiveComputeConstructInfo.Kind, GK, + E); + case OpenACCDirectiveKind::KernelsLoop: + case OpenACCDirectiveKind::Kernels: + return CheckGangKernelsExpr(*this, ExistingClauses, DK, + ActiveComputeConstructInfo.Kind, GK, E); + default: + llvm_unreachable("Non compute construct in active compute construct?"); } - } break; default: - llvm_unreachable("Non compute construct in active compute construct?"); + // TODO: OpenACC: when we implement this on 'routine', we'll have to + // implement its checking here. + llvm_unreachable("Invalid directive kind for a Gang clause"); } - llvm_unreachable("Compute construct directive not handled?"); } diff --git a/clang/lib/Sema/TreeTransform.h b/clang/lib/Sema/TreeTransform.h index 3a8f2d95f329b8..81e515e7cb2a9a 100644 --- a/clang/lib/Sema/TreeTransform.h +++ b/clang/lib/Sema/TreeTransform.h @@ -12027,7 +12027,9 @@ void OpenACCClauseTransform<Derived>::VisitGangClause( if (!ER.isUsable()) continue; - ER = Self.getSema().OpenACC().CheckGangExpr(C.getExpr(I).first, ER.get()); + ER = Self.getSema().OpenACC().CheckGangExpr(ExistingClauses, + ParsedClause.getDirectiveKind(), + C.getExpr(I).first, ER.get()); if (!ER.isUsable()) continue; TransformedGangKinds.push_back(C.getExpr(I).first); diff --git a/clang/test/AST/ast-print-openacc-combined-construct.cpp b/clang/test/AST/ast-print-openacc-combined-construct.cpp index 435c770c7457d1..40f174e539c01e 100644 --- a/clang/test/AST/ast-print-openacc-combined-construct.cpp +++ b/clang/test/AST/ast-print-openacc-combined-construct.cpp @@ -242,4 +242,75 @@ void foo() { #pragma acc parallel loop vector_length((int)array[1]) for(int i = 0;i<5;++i); +// CHECK: #pragma acc parallel loop gang(dim: 2) +// CHECK-NEXT: for (int i = 0; i < 5; ++i) +// CHECK-NEXT: ; +#pragma acc parallel loop gang(dim:2) + for(int i = 0;i<5;++i); + +// CHECK: #pragma acc serial loop gang(static: i) +// CHECK-NEXT: for (int i = 0; i < 5; ++i) +// CHECK-NEXT: ; +#pragma acc serial loop gang(static:i) + for(int i = 0;i<5;++i); + +// CHECK: #pragma acc parallel loop gang(static: i) gang(dim: 2) +// CHECK-NEXT: for (int i = 0; i < 5; ++i) +// CHECK-NEXT: ; +#pragma acc parallel loop gang(static:i) gang(dim:2) + for(int i = 0;i<5;++i); + +// CHECK: #pragma acc parallel loop gang(static: i, dim: 2) +// CHECK-NEXT: for (int i = 0; i < 5; ++i) +// CHECK-NEXT: ; +#pragma acc parallel loop gang(static:i, dim:2) + for(int i = 0;i<5;++i); + +// CHECK: #pragma acc parallel loop gang(dim: 2) +// CHECK-NEXT: for (int i = 0; i < 5; ++i) +// CHECK-NEXT: ; +#pragma acc parallel loop gang(dim:2) + for(int i = 0;i<5;++i); + +// CHECK: #pragma acc parallel loop gang(static: i) +// CHECK-NEXT: for (int i = 0; i < 5; ++i) +// CHECK-NEXT: ; +#pragma acc parallel loop gang(static:i) + for(int i = 0;i<5;++i); + +// CHECK: #pragma acc parallel loop gang(static: i) gang(dim: 2) +// CHECK-NEXT: for (int i = 0; i < 5; ++i) +// CHECK-NEXT: ; +#pragma acc parallel loop gang(static:i) gang(dim:2) + for(int i = 0;i<5;++i); + +// CHECK: #pragma acc parallel loop gang(static: i, dim: 2) +// CHECK-NEXT: for (int i = 0; i < 5; ++i) +// CHECK-NEXT: ; +#pragma acc parallel loop gang(static:i, dim:2) + for(int i = 0;i<5;++i); + +// CHECK: #pragma acc kernels loop gang(num: i) gang(static: i) +// CHECK-NEXT: for (int i = 0; i < 5; ++i) +// CHECK-NEXT: ; +#pragma acc kernels loop gang(i) gang(static:i) + for(int i = 0;i<5;++i); + +// CHECK: #pragma acc kernels loop gang(num: i) gang(static: i) +// CHECK-NEXT: for (int i = 0; i < 5; ++i) +// CHECK-NEXT: ; +#pragma acc kernels loop gang(num:i) gang(static:i) + for(int i = 0;i<5;++i); + +// CHECK: #pragma acc serial loop gang(static: i) +// CHECK-NEXT: for (int i = 0; i < 5; ++i) +// CHECK-NEXT: ; +#pragma acc serial loop gang(static:i) + for(int i = 0;i<5;++i); + +// CHECK: #pragma acc serial loop gang(static: *) +// CHECK-NEXT: for (int i = 0; i < 5; ++i) +// CHECK-NEXT: ; +#pragma acc serial loop gang(static:*) + for(int i = 0;i<5;++i); } diff --git a/clang/test/SemaOpenACC/combined-construct-auto_seq_independent-clauses.c b/clang/test/SemaOpenACC/combined-construct-auto_seq_independent-clauses.c index a6f57a63a91ddf..e35bd6da2f18b3 100644 --- a/clang/test/SemaOpenACC/combined-construct-auto_seq_independent-clauses.c +++ b/clang/test/SemaOpenACC/combined-construct-auto_seq_independent-clauses.c @@ -155,7 +155,6 @@ void uses() { #pragma acc parallel loop auto tile(1+2, 1) for(unsigned j = 0; j < 5; ++j) for(unsigned i = 0; i < 5; ++i); - // expected-warning@+1{{OpenACC clause 'gang' not yet implemented}} #pragma acc parallel loop auto gang for(unsigned i = 0; i < 5; ++i); #pragma acc parallel loop auto wait @@ -276,7 +275,6 @@ void uses() { #pragma acc parallel loop tile(1+2, 1) auto for(unsigned j = 0; j < 5; ++j) for(unsigned i = 0; i < 5; ++i); - // expected-warning@+1{{OpenACC clause 'gang' not yet implemented}} #pragma acc parallel loop gang auto for(unsigned i = 0; i < 5; ++i); #pragma acc parallel loop wait auto @@ -398,7 +396,6 @@ void uses() { #pragma acc parallel loop independent tile(1+2, 1) for(unsigned j = 0; j < 5; ++j) for(unsigned i = 0; i < 5; ++i); - // expected-warning@+1{{OpenACC clause 'gang' not yet implemented}} #pragma acc parallel loop independent gang for(unsigned i = 0; i < 5; ++i); #pragma acc parallel loop independent wait @@ -519,7 +516,6 @@ void uses() { #pragma acc parallel loop tile(1+2, 1) independent for(unsigned j = 0; j < 5; ++j) for(unsigned i = 0; i < 5; ++i); - // expected-warning@+1{{OpenACC clause 'gang' not yet implemented}} #pragma acc parallel loop gang independent for(unsigned i = 0; i < 5; ++i); #pragma acc parallel loop wait independent @@ -650,9 +646,8 @@ void uses() { #pragma acc parallel loop seq wait for(unsigned i = 0; i < 5; ++i); - // TODOexpected-error@+2{{OpenACC clause 'seq' may not appear on the same construct as a 'gang' clause on a 'parallel loop' construct}} - // TODOexpected-note@+1{{previous clause is here}} - // expected-warning@+1{{OpenACC clause 'gang' not yet implemented}} + // expected-error@+2{{OpenACC clause 'seq' may not appear on the same construct as a 'gang' clause on a 'parallel loop' construct}} + // expected-note@+1{{previous clause is here}} #pragma acc parallel loop gang seq for(unsigned i = 0; i < 5; ++i); // TODOexpected-error@+2{{OpenACC clause 'seq' may not appear on the same construct as a 'worker' clause on a 'parallel loop' construct}} diff --git a/clang/test/SemaOpenACC/combined-construct-device_type-clause.c b/clang/test/SemaOpenACC/combined-construct-device_type-clause.c index 9a60fb4c665e5a..d3ed8234d16b14 100644 --- a/clang/test/SemaOpenACC/combined-construct-device_type-clause.c +++ b/clang/test/SemaOpenACC/combined-construct-device_type-clause.c @@ -212,7 +212,6 @@ void uses() { for(int j = 0; j < 5; ++j) for(int i = 0; i < 5; ++i); - // expected-warning@+1{{OpenACC clause 'gang' not yet implemented, clause ignored}} #pragma acc serial loop dtype(*) gang for(int i = 0; i < 5; ++i); #pragma acc parallel loop device_type(*) wait diff --git a/clang/test/SemaOpenACC/combined-construct-gang-ast.cpp b/clang/test/SemaOpenACC/combined-construct-gang-ast.cpp new file mode 100644 index 00000000000000..f179b928215e71 --- /dev/null +++ b/clang/test/SemaOpenACC/combined-construct-gang-ast.cpp @@ -0,0 +1,216 @@ + +// RUN: %clang_cc1 %s -fopenacc -ast-dump | FileCheck %s + +// Test this with PCH. +// RUN: %clang_cc1 %s -fopenacc -emit-pch -o %t %s +// RUN: %clang_cc1 %s -fopenacc -include-pch %t -ast-dump-all | FileCheck %s +#ifndef PCH_HELPER +#define PCH_HELPER +void NormalUses() { + // CHECK: FunctionDecl{{.*}}NormalUses + // CHECK-NEXT: CompoundStmt + + int Val; + +#pragma acc parallel loop gang(dim:1) + for(int i = 0; i < 5; ++i); + // CHECK: OpenACCCombinedConstruct{{.*}} parallel loop + // CHECK-NEXT: gang clause dim + // CHECK-NEXT: ConstantExpr + // CHECK-NEXT: value: Int 1 + // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 1 + // CHECK-NEXT: ForStmt + // CHECK: NullStmt + +#pragma acc parallel loop gang(static:Val) + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop + // CHECK-NEXT: gang clause static + // CHECK-NEXT: ImplicitCastExpr + // CHECK-NEXT: DeclRefExpr{{.*}}'Val' 'int' + // CHECK-NEXT: ForStmt + // CHECK: NullStmt + +#pragma acc kernels loop gang(num:1) gang(static:Val) + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} kernels loop + // CHECK-NEXT: gang clause num + // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 1 + // CHECK-NEXT: gang clause static + // CHECK-NEXT: ImplicitCastExpr + // CHECK-NEXT: DeclRefExpr{{.*}}'Val' 'int' + // CHECK-NEXT: ForStmt + // CHECK: NullStmt + +#pragma acc parallel loop gang(dim:1, static:Val) + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop + // CHECK-NEXT: gang clause dim static + // CHECK-NEXT: ConstantExpr + // CHECK-NEXT: value: Int 1 + // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 1 + // CHECK-NEXT: ImplicitCastExpr + // CHECK-NEXT: DeclRefExpr{{.*}}'Val' 'int' + // CHECK-NEXT: ForStmt + // CHECK: NullStmt + +#pragma acc serial loop gang(static:Val) + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} serial loop + // CHECK-NEXT: gang clause static + // CHECK-NEXT: ImplicitCastExpr + // CHECK-NEXT: DeclRefExpr{{.*}}'Val' 'int' + // CHECK-NEXT: ForStmt + // CHECK: NullStmt + +#pragma acc serial loop gang(static:*) + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} serial loop + // CHECK-NEXT: gang clause static + // CHECK-NEXT: OpenACCAsteriskSizeExpr + // CHECK-NEXT: ForStmt + // CHECK: NullStmt + +#pragma acc serial loop gang + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} serial loop + // CHECK-NEXT: gang clause + // CHECK-NEXT: ForStmt + // CHECK: NullStmt + +#pragma acc kernels loop gang(num:1) + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} kernels loop + // CHECK-NEXT: gang clause num + // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 1 + // CHECK-NEXT: ForStmt + // CHECK: NullStmt +} + +template<typename T, unsigned One> +void TemplateUses(T Val) { + // CHECK: FunctionTemplateDecl{{.*}}TemplateUses + // CHECK-NEXT: TemplateTypeParmDecl {{.*}} referenced typename depth 0 index 0 T + // CHECK-NEXT: NonTypeTemplateParmDecl{{.*}} referenced 'unsigned int' depth 0 index 1 One + // CHECK-NEXT: FunctionDecl{{.*}} TemplateUses 'void (T)' + // CHECK-NEXT: ParmVarDecl{{.*}} referenced Val 'T' + // CHECK-NEXT: CompoundStmt + +#pragma acc parallel loop gang(dim:One) + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop + // CHECK-NEXT: gang clause dim + // CHECK-NEXT: DeclRefExpr{{.*}}'One' 'unsigned int' + // CHECK-NEXT: ForStmt + // CHECK: NullStmt + +#pragma acc parallel loop gang(static:Val) + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop + // CHECK-NEXT: gang clause static + // CHECK-NEXT: DeclRefExpr{{.*}}'Val' 'T' + // CHECK-NEXT: ForStmt + // CHECK: NullStmt + + +#pragma acc parallel loop gang(static:*) + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop + // CHECK-NEXT: gang clause static + // CHECK-NEXT: OpenACCAsteriskSizeExpr + // CHECK-NEXT: ForStmt + // CHECK: NullStmt + +#pragma acc parallel loop gang(dim:One) gang(static:Val) + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop + // CHECK-NEXT: gang clause dim + // CHECK-NEXT: DeclRefExpr{{.*}}'One' 'unsigned int' + // CHECK-NEXT: gang clause static + // CHECK-NEXT: DeclRefExpr{{.*}}'Val' 'T' + // CHECK-NEXT: ForStmt + // CHECK: NullStmt + +#pragma acc parallel loop gang(dim:One, static:Val) + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop + // CHECK-NEXT: gang clause dim static + // CHECK-NEXT: DeclRefExpr{{.*}}'One' 'unsigned int' + // CHECK-NEXT: DeclRefExpr{{.*}}'Val' 'T' + // CHECK-NEXT: ForStmt + // CHECK: NullStmt + +#pragma acc serial loop gang + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} serial loop + // CHECK-NEXT: gang clause + // CHECK-NEXT: ForStmt + // CHECK: NullStmt + + // Instantiation: + // CHECK-NEXT: FunctionDecl{{.*}} used TemplateUses 'void (int)' implicit_instantiation + // CHECK-NEXT: TemplateArgument type 'int' + // CHECK-NEXT: BuiltinType{{.*}} 'int' + // CHECK-NEXT: TemplateArgument integral '1U' + // CHECK-NEXT: ParmVarDecl{{.*}} used Val 'int' + // CHECK-NEXT: CompoundStmt + + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop + // CHECK-NEXT: gang clause dim + // CHECK-NEXT: ConstantExpr + // CHECK-NEXT: value: Int 1 + // CHECK-NEXT: SubstNonTypeTemplateParmExpr + // CHECK-NEXT: NonTypeTemplateParmDecl + // CHECK-NEXT: IntegerLiteral{{.*}}'unsigned int' 1 + // CHECK-NEXT: ForStmt + // CHECK: NullStmt + + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop + // CHECK-NEXT: gang clause static + // CHECK-NEXT: ImplicitCastExpr + // CHECK-NEXT: DeclRefExpr{{.*}}'Val' 'int' + // CHECK-NEXT: ForStmt + // CHECK: NullStmt + + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop + // CHECK-NEXT: gang clause static + // CHECK-NEXT: OpenACCAsteriskSizeExpr + // CHECK-NEXT: ForStmt + // CHECK: NullStmt + + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop + // CHECK-NEXT: gang clause dim + // CHECK-NEXT: ConstantExpr + // CHECK-NEXT: value: Int 1 + // CHECK-NEXT: SubstNonTypeTemplateParmExpr + // CHECK-NEXT: NonTypeTemplateParmDecl + // CHECK-NEXT: IntegerLiteral{{.*}}'unsigned int' 1 + // CHECK-NEXT: gang clause static + // CHECK-NEXT: ImplicitCastExpr + // CHECK-NEXT: DeclRefExpr{{.*}}'Val' 'int' + // CHECK-NEXT: ForStmt + // CHECK: NullStmt + + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop + // CHECK-NEXT: gang clause dim static + // CHECK-NEXT: ConstantExpr + // CHECK-NEXT: value: Int 1 + // CHECK-NEXT: SubstNonTypeTemplateParmExpr + // CHECK-NEXT: NonTypeTemplateParmDecl + // CHECK-NEXT: IntegerLiteral{{.*}}'unsigned int' 1 + // CHECK-NEXT: ImplicitCastExpr + // CHECK-NEXT: DeclRefExpr{{.*}}'Val' 'int' + // CHECK-NEXT: ForStmt + // CHECK: NullStmt + + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} serial loop + // CHECK-NEXT: gang clause + // CHECK-NEXT: ForStmt + // CHECK: NullStmt +} + +void inst() { + TemplateUses<int, 1>(5); +} + +#endif // PCH_HELPER diff --git a/clang/test/SemaOpenACC/combined-construct-gang-clause.cpp b/clang/test/SemaOpenACC/combined-construct-gang-clause.cpp new file mode 100644 index 00000000000000..c5346c4cf5b90a --- /dev/null +++ b/clang/test/SemaOpenACC/combined-construct-gang-clause.cpp @@ -0,0 +1,307 @@ +// RUN: %clang_cc1 %s -fopenacc -verify + +struct S{}; +struct Converts{ + operator int(); +}; + +template<typename T, unsigned Zero, unsigned Two, unsigned Four> +void ParallelTempl() { + T i; + + // expected-error@+1{{'num' argument on 'gang' clause is not permitted on a 'parallel loop' construct}} +#pragma acc parallel loop gang(i) + for(int i = 0; i < 5; ++i); + + // expected-error@+1{{'num' argument on 'gang' clause is not permitted on a 'parallel loop' construct}} +#pragma acc parallel loop gang(num:i) + for(int i = 0; i < 5; ++i); + + // expected-error@+1{{argument to 'gang' clause dimension must be a constant expression}} +#pragma acc parallel loop gang(dim:i) + for(int i = 0; i < 5; ++i); + + // expected-error@+1{{argument to 'gang' clause dimension must be 1, 2, or 3: evaluated to 0}} +#pragma acc parallel loop gang(dim:Zero) + for(int i = 0; i < 5; ++i); + +#pragma acc parallel loop gang(dim:Two) + for(int i = 0; i < 5; ++i); + + // expected-error@+1{{argument to 'gang' clause dimension must be 1, 2, or 3: evaluated to 4}} +#pragma acc parallel loop gang(dim:Four) + for(int i = 0; i < 5; ++i); + +#pragma acc parallel loop gang(static:i) gang(dim:Two) + for(int i = 0; i < 5; ++i); + +#pragma acc parallel loop gang(static:i, dim:Two) + for(int i = 0; i < 5; ++i); + +#pragma acc parallel loop gang(dim:Two) gang(static:*) + for(int i = 0; i < 5; ++i); + +#pragma acc parallel loop gang(dim:Two, static:*) + for(int i = 0; i < 5; ++i); + + // expected-error@+4{{OpenACC 'gang' clause may have at most one 'static' argument}} + // expected-note@+3{{previous expression is here}} + // expected-error@+2{{OpenACC 'gang' clause may have at most one 'dim' argument}} + // expected-note@+1{{previous expression is here}} +#pragma acc parallel loop gang(dim:Two, static:*, dim:1, static:i) + for(int i = 0; i < 5; ++i); +} + +void Parallel() { + ParallelTempl<int, 0, 2, 4>(); // expected-note{{in instantiation of function template}} + int i; + + // expected-error@+1{{'num' argument on 'gang' clause is not permitted on a 'parallel loop' construct}} +#pragma acc parallel loop gang(i) + for(int i = 0; i < 5; ++i); + + // expected-error@+1{{'num' argument on 'gang' clause is not permitted on a 'parallel loop' construct}} +#pragma acc parallel loop gang(num:i) + for(int i = 0; i < 5; ++i); + + // expected-error@+1{{argument to 'gang' clause dimension must be a constant expression}} +#pragma acc parallel loop gang(dim:i) + for(int i = 0; i < 5; ++i); + + // expected-error@+1{{argument to 'gang' clause dimension must be 1, 2, or 3: evaluated to 0}} +#pragma acc parallel loop gang(dim:0) + for(int i = 0; i < 5; ++i); + +#pragma acc parallel loop gang(dim:2) + for(int i = 0; i < 5; ++i); + + // expected-error@+1{{argument to 'gang' clause dimension must be 1, 2, or 3: evaluated to 4}} +#pragma acc parallel loop gang(dim:4) + for(int i = 0; i < 5; ++i); + +#pragma acc parallel loop gang(static:i) gang(dim:2) + for(int i = 0; i < 5; ++i); + +#pragma acc parallel loop gang(static:i, dim:2) + for(int i = 0; i < 5; ++i); + +#pragma acc parallel loop gang(dim:2) gang(static:*) + for(int i = 0; i < 5; ++i); + +#pragma acc parallel loop gang(dim:2, static:*) + for(int i = 0; i < 5; ++i); + + // expected-error@+4{{OpenACC 'gang' clause may have at most one 'static' argument}} + // expected-note@+3{{previous expression is here}} + // expected-error@+2{{OpenACC 'gang' clause may have at most one 'dim' argument}} + // expected-note@+1{{previous expression is here}} +#pragma acc parallel loop gang(dim:2, static:*, dim:1, static:i) + for(int i = 0; i < 5; ++i); +} + +template<typename SomeS, typename SomeC, typename Int> +void StaticIsIntegralTempl() { + SomeS s; + // expected-error@+1{{OpenACC clause 'gang' requires expression of integer type ('S' invalid)}} +#pragma acc parallel loop gang(dim:2) gang(static:s) + for(int i = 0; i < 5; ++i); + + SomeC C; +#pragma acc parallel loop gang(dim:2) gang(static:C) + for(int i = 0; i < 5; ++i); + Int i; +#pragma acc parallel loop gang(dim:2) gang(static:i) + for(int i = 0; i < 5; ++i); + +#pragma acc parallel loop gang(dim:2) gang(static:*) + for(int i = 0; i < 5; ++i); +} + +void StaticIsIntegral() { + StaticIsIntegralTempl<S, Converts, int>();// expected-note{{in instantiation of function template}} + + S s; + // expected-error@+1{{OpenACC clause 'gang' requires expression of integer type ('S' invalid)}} +#pragma acc parallel loop gang(dim:2) gang(static:s) + for(int i = 0; i < 5; ++i); + + Converts C; +#pragma acc parallel loop gang(dim:2) gang(static:C) + for(int i = 0; i < 5; ++i); +} + +template<unsigned I> +void SerialTempl() { + // expected-error@+1{{'num' argument on 'gang' clause is not permitted on a 'serial loop'}} +#pragma acc serial loop gang(I) + for(int i = 0; i < 5; ++i); + + // expected-error@+1{{'num' argument on 'gang' clause is not permitted on a 'serial loop'}} +#pragma acc serial loop gang(num:I) + for(int i = 0; i < 5; ++i); + + // expected-error@+1{{'dim' argument on 'gang' clause is not permitted on a 'serial loop'}} +#pragma acc serial loop gang(dim:I) + for(int i = 0; i < 5; ++i); + +#pragma acc serial loop gang(static:I) + for(int i = 0; i < 5; ++i); +} + +void Serial() { + SerialTempl<2>(); + + // expected-error@+1{{'num' argument on 'gang' clause is not permitted on a 'serial loop'}} +#pragma acc serial loop gang(1) + for(int i = 0; i < 5; ++i); + + // expected-error@+1{{'num' argument on 'gang' clause is not permitted on a 'serial loop'}} +#pragma acc serial loop gang(num:1) + for(int i = 0; i < 5; ++i); + + // expected-error@+1{{'dim' argument on 'gang' clause is not permitted on a 'serial loop'}} +#pragma acc serial loop gang(dim:1) + for(int i = 0; i < 5; ++i); + +#pragma acc serial loop gang(static:1) + for(int i = 0; i < 5; ++i); + + int i; +#pragma acc serial loop gang(static:i) + for(int i = 0; i < 5; ++i); +} + +template<typename T> +void KernelsTempl() { + T t; + // expected-error@+1{{'dim' argument on 'gang' clause is not permitted on a 'kernels loop'}} +#pragma acc kernels loop gang(dim:t) + for(int i = 0; i < 5; ++i); + +#pragma acc kernels loop gang(static:t) + for(int i = 0; i < 5; ++i); + + // expected-error@+2{{'num_gangs' clause not allowed on a 'kernels loop' construct that has a 'gang' clause with a 'num' argument}} + // expected-note@+1{{previous clause is here}} +#pragma acc kernels loop gang(t) num_gangs(t) + for(int i = 0; i < 5; ++i); + + // OK, kernels loop should block this. +#pragma acc kernels num_gangs(t) +#pragma acc kernels loop gang(static:t) + for(int i = 0; i < 5; ++i); + + // expected-error@+2{{'num' argument to 'gang' clause not allowed on a 'kernels loop' construct that has a 'num_gangs' clause}} + // expected-note@+1{{previous clause is here}} +#pragma acc kernels loop num_gangs(t) gang(t) + for(int i = 0; i < 5; ++i); + + // expected-error@+2{{'num_gangs' clause not allowed on a 'kernels loop' construct that has a 'gang' clause with a 'num' argument}} + // expected-note@+1{{previous clause is here}} +#pragma acc kernels loop gang(num:t) num_gangs(t) + for(int i = 0; i < 5; ++i); + + // expected-error@+2{{'num' argument to 'gang' clause not allowed on a 'kernels loop' construct that has a 'num_gangs' clause}} + // expected-note@+1{{previous clause is here}} +#pragma acc kernels loop num_gangs(t) gang(num:t) + for(int i = 0; i < 5; ++i); +} + +void Kernels() { + KernelsTempl<unsigned>(); + + // expected-error@+1{{'dim' argument on 'gang' clause is not permitted on a 'kernels loop'}} +#pragma acc kernels loop gang(dim:1) + for(int i = 0; i < 5; ++i); + + unsigned t; +#pragma acc kernels loop gang(static:t) + for(int i = 0; i < 5; ++i); + + // expected-error@+2{{'num_gangs' clause not allowed on a 'kernels loop' construct that has a 'gang' clause with a 'num' argument}} + // expected-note@+1{{previous clause is here}} +#pragma acc kernels loop gang(t) num_gangs(t) + for(int i = 0; i < 5; ++i); + + // expected-error@+2{{'num' argument to 'gang' clause not allowed on a 'kernels loop' construct that has a 'num_gangs' clause}} + // expected-note@+1{{previous clause is here}} +#pragma acc kernels loop num_gangs(t) gang(t) + for(int i = 0; i < 5; ++i); + + // expected-error@+2{{'num_gangs' clause not allowed on a 'kernels loop' construct that has a 'gang' clause with a 'num' argument}} + // expected-note@+1{{previous clause is here}} +#pragma acc kernels loop gang(num:t) num_gangs(t) + for(int i = 0; i < 5; ++i); + + // expected-error@+2{{'num' argument to 'gang' clause not allowed on a 'kernels loop' construct that has a 'num_gangs' clause}} + // expected-note@+1{{previous clause is here}} +#pragma acc kernels loop num_gangs(t) gang(num:t) + for(int i = 0; i < 5; ++i); + + // OK, intervening compute/combined construct. +#pragma acc kernels loop gang(num:1) + for(int i = 0; i < 5; ++i) { +#pragma acc serial loop gang(static:1) + for(int i = 0; i < 5; ++i); + } +#pragma acc kernels loop gang(num:1) + for(int i = 0; i < 5; ++i) { +#pragma acc serial +#pragma acc loop gang(static:1) + for(int i = 0; i < 5; ++i); + } +#pragma acc kernels +#pragma acc loop gang(num:1) + for(int i = 0; i < 5; ++i) { +#pragma acc serial loop gang(static:1) + for(int i = 0; i < 5; ++i); + } + +#pragma acc kernels loop gang(num:1) + for(int i = 0; i < 5; ++i) { + // expected-error@+2{{loop with a 'gang' clause may not exist in the region of a 'gang' clause on a 'kernels loop' construct}} + // expected-note@-3{{previous clause is here}} +#pragma acc loop gang(static:1) + for(int i = 0; i < 5; ++i); + } + + // OK, on a diff erent 'loop', not in the associated stmt. +#pragma acc kernels loop gang(num:1) + for(int i = 0; i < 5; ++i); +#pragma acc loop gang(static:1) + for(int i = 0; i < 5; ++i); + + // OK, on a diff erent 'loop', not in the associated stmt. +#pragma acc kernels loop gang(num:1) + for(int i = 0; i < 5; ++i); +#pragma acc kernels loop gang(static:1) + for(int i = 0; i < 5; ++i); + + // expected-error@+2{{OpenACC 'gang' clause may have at most one unnamed or 'num' argument}} + // expected-note@+1{{previous expression is here}} +#pragma acc kernels loop gang(5, num:1) + for(int i = 0; i < 5; ++i); + + // expected-error@+2{{OpenACC 'gang' clause may have at most one unnamed or 'num' argument}} + // expected-note@+1{{previous expression is here}} +#pragma acc kernels loop gang(num:5, 1) + for(int i = 0; i < 5; ++i); + + // expected-error@+3{{OpenACC 'gang' clause may have at most one unnamed or 'num' argument}} + // expected-note@+2{{previous expression is here}} +#pragma acc kernels +#pragma acc kernels loop gang(num:5, num:1) + for(int i = 0; i < 5; ++i); +} + +void MaxOneEntry() { + // expected-error@+2{{OpenACC 'gang' clause may have at most one 'static' argument}} + // expected-note@+1{{previous expression is here}} +#pragma acc kernels loop gang(static: 1, static:1) + for(int i = 0; i < 5; ++i); + +#pragma acc kernels loop gang gang(static:1) + for(int i = 0; i < 5; ++i); +} + + diff --git a/clang/test/SemaOpenACC/loop-construct-gang-clause.cpp b/clang/test/SemaOpenACC/loop-construct-gang-clause.cpp index a2bda5a7e82f63..6b4c03bbc50c6b 100644 --- a/clang/test/SemaOpenACC/loop-construct-gang-clause.cpp +++ b/clang/test/SemaOpenACC/loop-construct-gang-clause.cpp @@ -278,7 +278,7 @@ void Kernels() { #pragma acc kernels #pragma acc loop gang(num:1) for(int j = 0; j < 5; ++j) { - // expected-error@+2{{loop with a 'gang' clause may not exist in the region of a 'gang' clause on a 'kernels' compute construct}} + // expected-error@+2{{loop with a 'gang' clause may not exist in the region of a 'gang' clause on a 'kernels' construct}} // expected-note@-3{{previous clause is here}} #pragma acc loop gang(static:1) for(int i = 0; i < 5; ++i); _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
