Skip to content

Commit

Permalink
[OpenACC] 'vector' clause implementation for combined constructs
Browse files Browse the repository at this point in the history 
Similar to 'worker', the 'vector' clause has some rules that needed to
be applied on its argument legality that for combined constructs need to
look at the current construct, not the 'effective' parent construct.
Additionally, it has some interaction with `vector_length` that needed
to be encoded as well.  This patch implements it.
  • Loading branch information
@erichkeane
erichkeane committed Dec 6, 2024
1 parent f09b16e commit 009b5e8
Show file tree
Hide file tree
Showing 8 changed files with 507 additions and 65 deletions.
6 changes: 3 additions & 3 deletions clang/include/clang/Basic/DiagnosticSemaKinds.td
View file
Original file line number Diff line number Diff line change
Expand Up @@ -12762,9 +12762,9 @@ 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 '%1' "
"construct%select{| associated with a '%3' construct}2 that has a "
"'%4' clause">;
: Error<"'%0' argument to '%1' clause not allowed on a '%2' "
"construct%select{| associated with a '%4' construct}3 that has a "
"'%5' clause">;
def err_acc_num_arg_conflict_reverse
: Error<"'%0' clause not allowed on a 'kernels loop' construct that "
"has a '%1' clause with a%select{n| 'num'}2 argument">;
Expand Down
159 changes: 109 additions & 50 deletions clang/lib/Sema/SemaOpenACC.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -804,6 +804,25 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitVectorLengthClause(
if (checkAlreadyHasClauseOfKind(SemaRef, ExistingClauses, Clause))
return nullptr;

// OpenACC 3.3 Section 2.9.4:
// An argument is allowed only when the 'vector_length' does not appear on the
// 'kernels' construct.
if (Clause.getDirectiveKind() == OpenACCDirectiveKind::KernelsLoop) {
auto VectorClauses = llvm::make_filter_range(
ExistingClauses, llvm::IsaPred<OpenACCVectorClause>);

for (auto *VC : VectorClauses) {
if (cast<OpenACCVectorClause>(VC)->hasIntExpr()) {
SemaRef.Diag(Clause.getBeginLoc(),
diag::err_acc_num_arg_conflict_reverse)
<< OpenACCClauseKind::VectorLength << OpenACCClauseKind::Vector
<< /*num argument*/ 0;
SemaRef.Diag(VC->getBeginLoc(), diag::note_acc_previous_clause_here);
return nullptr;
}
}
}

assert(Clause.getIntExprs().size() == 1 &&
"Invalid number of expressions for NumWorkers");
return OpenACCVectorLengthClause::Create(
Expand Down Expand Up @@ -1097,6 +1116,14 @@ ExprResult DiagIntArgInvalid(SemaOpenACC &S, Expr *E, OpenACCGangKind GK,
<< HasAssocKind(DK, AssocKind) << AssocKind;
return ExprError();
}
ExprResult DiagIntArgInvalid(SemaOpenACC &S, Expr *E, StringRef TagKind,
OpenACCClauseKind CK, OpenACCDirectiveKind DK,
OpenACCDirectiveKind AssocKind) {
S.Diag(E->getBeginLoc(), diag::err_acc_int_arg_invalid)
<< TagKind << CK << IsOrphanLoop(DK, AssocKind) << DK
<< HasAssocKind(DK, AssocKind) << AssocKind;
return ExprError();
}

ExprResult CheckGangParallelExpr(SemaOpenACC &S, OpenACCDirectiveKind DK,
OpenACCDirectiveKind AssocKind,
Expand Down Expand Up @@ -1172,8 +1199,9 @@ ExprResult CheckGangKernelsExpr(SemaOpenACC &S,

if (Itr != Collection.end()) {
S.Diag(E->getBeginLoc(), diag::err_acc_num_arg_conflict)
<< OpenACCClauseKind::Gang << DK << HasAssocKind(DK, AssocKind)
<< AssocKind << OpenACCClauseKind::NumGangs;
<< "num" << OpenACCClauseKind::Gang << DK
<< HasAssocKind(DK, AssocKind) << AssocKind
<< OpenACCClauseKind::NumGangs;

S.Diag((*Itr)->getBeginLoc(), diag::note_acc_previous_clause_here);
return ExprError();
Expand Down Expand Up @@ -1206,63 +1234,94 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitVectorClause(
SemaOpenACC::OpenACCParsedClause &Clause) {
if (DiagIfSeqClause(Clause))
return nullptr;
// 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)

// Restrictions only properly implemented on 'loop'/'combined' constructs, and
// it is the only construct that can do anything with this, so skip/treat as
// unimplemented for the routine constructs.
if (Clause.getDirectiveKind() != OpenACCDirectiveKind::Loop &&
!isOpenACCCombinedDirectiveKind(Clause.getDirectiveKind()))
return isNotImplemented();

Expr *IntExpr =
Clause.getNumIntExprs() != 0 ? Clause.getIntExprs()[0] : nullptr;
if (IntExpr) {
switch (SemaRef.getActiveComputeConstructInfo().Kind) {
case OpenACCDirectiveKind::Invalid:
case OpenACCDirectiveKind::Parallel:
// No restriction on when 'parallel' can contain an argument.
break;
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.
DiagIntArgInvalid(SemaRef, IntExpr, OpenACCGangKind::Num,
OpenACCClauseKind::Vector, Clause.getDirectiveKind(),
SemaRef.getActiveComputeConstructInfo().Kind);
IntExpr = nullptr;
break;
case OpenACCDirectiveKind::Kernels: {
const auto *Itr =
llvm::find_if(SemaRef.getActiveComputeConstructInfo().Clauses,
llvm::IsaPred<OpenACCVectorLengthClause>);
if (Itr != SemaRef.getActiveComputeConstructInfo().Clauses.end()) {
SemaRef.Diag(IntExpr->getBeginLoc(), diag::err_acc_num_arg_conflict)
<< 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);
if (!isOpenACCCombinedDirectiveKind(Clause.getDirectiveKind())) {
switch (SemaRef.getActiveComputeConstructInfo().Kind) {
case OpenACCDirectiveKind::Invalid:
case OpenACCDirectiveKind::Parallel:
// No restriction on when 'parallel' can contain an argument.
break;
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.
DiagIntArgInvalid(SemaRef, IntExpr, "length", OpenACCClauseKind::Vector,
Clause.getDirectiveKind(),
SemaRef.getActiveComputeConstructInfo().Kind);
IntExpr = nullptr;
break;
case OpenACCDirectiveKind::Kernels: {
const auto *Itr =
llvm::find_if(SemaRef.getActiveComputeConstructInfo().Clauses,
llvm::IsaPred<OpenACCVectorLengthClause>);
if (Itr != SemaRef.getActiveComputeConstructInfo().Clauses.end()) {
SemaRef.Diag(IntExpr->getBeginLoc(), diag::err_acc_num_arg_conflict)
<< "length" << 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);

IntExpr = nullptr;
}
break;
}
default:
llvm_unreachable("Non compute construct in active compute construct");
}
} else {
if (Clause.getDirectiveKind() == OpenACCDirectiveKind::SerialLoop) {
DiagIntArgInvalid(SemaRef, IntExpr, "length", OpenACCClauseKind::Vector,
Clause.getDirectiveKind(),
SemaRef.getActiveComputeConstructInfo().Kind);
IntExpr = nullptr;
} else if (Clause.getDirectiveKind() ==
OpenACCDirectiveKind::KernelsLoop) {
const auto *Itr = llvm::find_if(
ExistingClauses, llvm::IsaPred<OpenACCVectorLengthClause>);
if (Itr != ExistingClauses.end()) {
SemaRef.Diag(IntExpr->getBeginLoc(), diag::err_acc_num_arg_conflict)
<< "length" << 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);

IntExpr = nullptr;
}
}
break;
}
default:
llvm_unreachable("Non compute construct in active compute construct");
}
}

// 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::Vector << OpenACCClauseKind::Vector
<< /*skip kernels construct info*/ 0;
SemaRef.Diag(SemaRef.LoopVectorClauseLoc,
diag::note_acc_previous_clause_here);
return nullptr;
if (!isOpenACCCombinedDirectiveKind(Clause.getDirectiveKind())) {
// 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::Vector << OpenACCClauseKind::Vector
<< /*skip kernels construct info*/ 0;
SemaRef.Diag(SemaRef.LoopVectorClauseLoc,
diag::note_acc_previous_clause_here);
return nullptr;
}
}

return OpenACCVectorClause::Create(Ctx, Clause.getBeginLoc(),
Expand Down Expand Up @@ -1305,7 +1364,7 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitWorkerClause(
llvm::IsaPred<OpenACCNumWorkersClause>);
if (Itr != SemaRef.getActiveComputeConstructInfo().Clauses.end()) {
SemaRef.Diag(IntExpr->getBeginLoc(), diag::err_acc_num_arg_conflict)
<< OpenACCClauseKind::Worker << Clause.getDirectiveKind()
<< "num" << OpenACCClauseKind::Worker << Clause.getDirectiveKind()
<< HasAssocKind(Clause.getDirectiveKind(),
SemaRef.getActiveComputeConstructInfo().Kind)
<< SemaRef.getActiveComputeConstructInfo().Kind
Expand Down Expand Up @@ -1334,7 +1393,7 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitWorkerClause(
llvm::IsaPred<OpenACCNumWorkersClause>);
if (Itr != ExistingClauses.end()) {
SemaRef.Diag(IntExpr->getBeginLoc(), diag::err_acc_num_arg_conflict)
<< OpenACCClauseKind::Worker << Clause.getDirectiveKind()
<< "num" << OpenACCClauseKind::Worker << Clause.getDirectiveKind()
<< HasAssocKind(Clause.getDirectiveKind(),
SemaRef.getActiveComputeConstructInfo().Kind)
<< SemaRef.getActiveComputeConstructInfo().Kind
Expand Down
42 changes: 42 additions & 0 deletions clang/test/AST/ast-print-openacc-combined-construct.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -344,4 +344,46 @@ void foo() {
#pragma acc kernels loop worker(num:5)
for(int i = 0;i<5;++i);

// CHECK: #pragma acc parallel loop vector
// CHECK-NEXT: for (int i = 0; i < 5; ++i)
// CHECK-NEXT: ;
#pragma acc parallel loop vector
for(int i = 0;i<5;++i);

// CHECK: #pragma acc parallel loop vector(length: 5)
// CHECK-NEXT: for (int i = 0; i < 5; ++i)
// CHECK-NEXT: ;
#pragma acc parallel loop vector(5)
for(int i = 0;i<5;++i);

// CHECK: #pragma acc parallel loop vector(length: 5)
// CHECK-NEXT: for (int i = 0; i < 5; ++i)
// CHECK-NEXT: ;
#pragma acc parallel loop vector(length:5)
for(int i = 0;i<5;++i);

// CHECK-NEXT: #pragma acc kernels loop vector
// CHECK-NEXT: for (int i = 0; i < 5; ++i)
// CHECK-NEXT: ;
#pragma acc kernels loop vector
for(int i = 0;i<5;++i);

// CHECK-NEXT: #pragma acc kernels loop vector(length: 5)
// CHECK-NEXT: for (int i = 0; i < 5; ++i)
// CHECK-NEXT: ;
#pragma acc kernels loop vector(5)
for(int i = 0;i<5;++i);

// CHECK-NEXT: #pragma acc kernels loop vector(length: 5)
// CHECK-NEXT: for (int i = 0; i < 5; ++i)
// CHECK-NEXT: ;
#pragma acc kernels loop vector(length:5)
for(int i = 0;i<5;++i);

// CHECK-NEXT: #pragma acc serial loop vector
// CHECK-NEXT: for (int i = 0; i < 5; ++i)
// CHECK-NEXT: ;
#pragma acc serial loop vector
for(int i = 0;i<5;++i);

}
9 changes: 2 additions & 7 deletions clang/test/SemaOpenACC/combined-construct-auto_seq_independent-clauses.c
View file
Original file line number Diff line number Diff line change
Expand Up @@ -48,7 +48,6 @@ void uses() {
for(unsigned i = 0; i < 5; ++i);
#pragma acc parallel loop auto worker
for(unsigned i = 0; i < 5; ++i);
// expected-warning@+1{{OpenACC clause 'vector' not yet implemented}}
#pragma acc parallel loop auto vector
for(unsigned i = 0; i < 5; ++i);
// expected-warning@+1{{OpenACC clause 'nohost' not yet implemented}}
Expand Down Expand Up @@ -167,7 +166,6 @@ void uses() {
for(unsigned i = 0; i < 5; ++i);
#pragma acc parallel loop worker auto
for(unsigned i = 0; i < 5; ++i);
// expected-warning@+1{{OpenACC clause 'vector' not yet implemented}}
#pragma acc parallel loop vector auto
for(unsigned i = 0; i < 5; ++i);
// expected-warning@+1{{OpenACC clause 'nohost' not yet implemented}}
Expand Down Expand Up @@ -287,7 +285,6 @@ void uses() {
for(unsigned i = 0; i < 5; ++i);
#pragma acc parallel loop independent worker
for(unsigned i = 0; i < 5; ++i);
// expected-warning@+1{{OpenACC clause 'vector' not yet implemented}}
#pragma acc parallel loop independent vector
for(unsigned i = 0; i < 5; ++i);
// expected-warning@+1{{OpenACC clause 'nohost' not yet implemented}}
Expand Down Expand Up @@ -406,7 +403,6 @@ void uses() {
for(unsigned i = 0; i < 5; ++i);
#pragma acc parallel loop worker independent
for(unsigned i = 0; i < 5; ++i);
// expected-warning@+1{{OpenACC clause 'vector' not yet implemented}}
#pragma acc parallel loop vector independent
for(unsigned i = 0; i < 5; ++i);
// expected-warning@+1{{OpenACC clause 'nohost' not yet implemented}}
Expand Down Expand Up @@ -650,9 +646,8 @@ void uses() {
// expected-note@+1{{previous clause is here}}
#pragma acc parallel loop worker seq
for(unsigned i = 0; i < 5; ++i);
// TODOexpected-error@+2{{OpenACC clause 'seq' may not appear on the same construct as a 'vector' clause on a 'parallel loop' construct}}
// TODOexpected-note@+1{{previous clause is here}}
// expected-warning@+1{{OpenACC clause 'vector' not yet implemented}}
// expected-error@+2{{OpenACC clause 'seq' may not appear on the same construct as a 'vector' clause on a 'parallel loop' construct}}
// expected-note@+1{{previous clause is here}}
#pragma acc parallel loop vector seq
for(unsigned i = 0; i < 5; ++i);
// expected-warning@+1{{OpenACC clause 'finalize' not yet implemented}}
Expand Down
1 change: 0 additions & 1 deletion clang/test/SemaOpenACC/combined-construct-device_type-clause.c
View file
Original file line number Diff line number Diff line change
Expand Up @@ -39,7 +39,6 @@ void uses() {
// 'worker', 'vector', 'seq', 'independent', 'auto', and 'tile' after
// 'device_type'.

//expected-warning@+1{{OpenACC clause 'vector' not yet implemented, clause ignored}}
#pragma acc parallel loop device_type(*) vector
for(int i = 0; i < 5; ++i);

Expand Down
Loading

0 comments on commit 009b5e8

Please sign in to comment.