Skip to content

Commit b8adf16

Browse files
erichkeaneerichkeane
committed
[OpenACC] Implement 'vector_length' clause On compute constructs
The 'vector_length' clause is semantically identical to the 'num_workers' clause, in that it takes a mandatory single int-expr. This is implemented identically to it.
1 parent 67669ea commit b8adf16

17 files changed

+328
-12
lines changed

‎clang/include/clang/AST/OpenACCClause.h

Lines changed: 10 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -196,6 +196,16 @@ class OpenACCNumWorkersClause : public OpenACCClauseWithSingleIntExpr {
196196
Expr *IntExpr, SourceLocation EndLoc);
197197
};
198198

199+
class OpenACCVectorLengthClause : public OpenACCClauseWithSingleIntExpr {
200+
OpenACCVectorLengthClause(SourceLocation BeginLoc, SourceLocation LParenLoc,
201+
Expr *IntExpr, SourceLocation EndLoc);
202+
203+
public:
204+
static OpenACCVectorLengthClause *
205+
Create(const ASTContext &C, SourceLocation BeginLoc, SourceLocation LParenLoc,
206+
Expr *IntExpr, SourceLocation EndLoc);
207+
};
208+
199209
template <class Impl> class OpenACCClauseVisitor {
200210
Impl &getDerived() { return static_cast<Impl &>(*this); }
201211

‎clang/include/clang/Basic/OpenACCClauses.def

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -19,5 +19,6 @@ VISIT_CLAUSE(Default)
1919
VISIT_CLAUSE(If)
2020
VISIT_CLAUSE(Self)
2121
VISIT_CLAUSE(NumWorkers)
22+
VISIT_CLAUSE(VectorLength)
2223

2324
#undef VISIT_CLAUSE

‎clang/include/clang/Sema/SemaOpenACC.h

Lines changed: 6 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -93,13 +93,15 @@ class SemaOpenACC : public SemaBase {
9393
}
9494

9595
unsigned getNumIntExprs() const {
96-
assert(ClauseKind == OpenACCClauseKind::NumWorkers &&
96+
assert((ClauseKind == OpenACCClauseKind::NumWorkers ||
97+
ClauseKind == OpenACCClauseKind::VectorLength) &&
9798
"Parsed clause kind does not have a int exprs");
9899
return std::get<IntExprDetails>(Details).IntExprs.size();
99100
}
100101

101102
ArrayRef<Expr *> getIntExprs() {
102-
assert(ClauseKind == OpenACCClauseKind::NumWorkers &&
103+
assert((ClauseKind == OpenACCClauseKind::NumWorkers ||
104+
ClauseKind == OpenACCClauseKind::VectorLength) &&
103105
"Parsed clause kind does not have a int exprs");
104106
return std::get<IntExprDetails>(Details).IntExprs;
105107
}
@@ -132,7 +134,8 @@ class SemaOpenACC : public SemaBase {
132134
}
133135

134136
void setIntExprDetails(ArrayRef<Expr *> IntExprs) {
135-
assert(ClauseKind == OpenACCClauseKind::NumWorkers &&
137+
assert((ClauseKind == OpenACCClauseKind::NumWorkers ||
138+
ClauseKind == OpenACCClauseKind::VectorLength) &&
136139
"Parsed clause kind does not have a int exprs");
137140
Details = IntExprDetails{{IntExprs.begin(), IntExprs.end()}};
138141
}

‎clang/lib/AST/OpenACCClause.cpp

Lines changed: 26 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -103,6 +103,27 @@ OpenACCNumWorkersClause::Create(const ASTContext &C, SourceLocation BeginLoc,
103103
OpenACCNumWorkersClause(BeginLoc, LParenLoc, IntExpr, EndLoc);
104104
}
105105

106+
OpenACCVectorLengthClause::OpenACCVectorLengthClause(SourceLocation BeginLoc,
107+
SourceLocation LParenLoc,
108+
Expr *IntExpr,
109+
SourceLocation EndLoc)
110+
: OpenACCClauseWithSingleIntExpr(OpenACCClauseKind::VectorLength, BeginLoc,
111+
LParenLoc, IntExpr, EndLoc) {
112+
assert((!IntExpr || IntExpr->isInstantiationDependent() ||
113+
IntExpr->getType()->isIntegerType()) &&
114+
"Condition expression type not scalar/dependent");
115+
}
116+
117+
OpenACCVectorLengthClause *
118+
OpenACCVectorLengthClause::Create(const ASTContext &C, SourceLocation BeginLoc,
119+
SourceLocation LParenLoc, Expr *IntExpr,
120+
SourceLocation EndLoc) {
121+
void *Mem = C.Allocate(sizeof(OpenACCVectorLengthClause),
122+
alignof(OpenACCVectorLengthClause));
123+
return new (Mem)
124+
OpenACCVectorLengthClause(BeginLoc, LParenLoc, IntExpr, EndLoc);
125+
}
126+
106127
//===----------------------------------------------------------------------===//
107128
// OpenACC clauses printing methods
108129
//===----------------------------------------------------------------------===//
@@ -124,3 +145,8 @@ void OpenACCClausePrinter::VisitNumWorkersClause(
124145
const OpenACCNumWorkersClause &C) {
125146
OS << "num_workers(" << C.getIntExpr() << ")";
126147
}
148+
149+
void OpenACCClausePrinter::VisitVectorLengthClause(
150+
const OpenACCVectorLengthClause &C) {
151+
OS << "vector_length(" << C.getIntExpr() << ")";
152+
}

‎clang/lib/AST/StmtProfile.cpp

Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -2503,6 +2503,12 @@ void OpenACCClauseProfiler::VisitNumWorkersClause(
25032503
Profiler.VisitStmt(Clause.getIntExpr());
25042504
}
25052505

2506+
void OpenACCClauseProfiler::VisitVectorLengthClause(
2507+
const OpenACCVectorLengthClause &Clause) {
2508+
assert(Clause.hasIntExpr() &&
2509+
"vector_length clause requires a valid int expr");
2510+
Profiler.VisitStmt(Clause.getIntExpr());
2511+
}
25062512
} // namespace
25072513

25082514
void StmtProfiler::VisitOpenACCComputeConstruct(

‎clang/lib/AST/TextNodeDumper.cpp

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -400,6 +400,7 @@ void TextNodeDumper::Visit(const OpenACCClause *C) {
400400
case OpenACCClauseKind::If:
401401
case OpenACCClauseKind::Self:
402402
case OpenACCClauseKind::NumWorkers:
403+
case OpenACCClauseKind::VectorLength:
403404
// The condition expression will be printed as a part of the 'children',
404405
// but print 'clause' here so it is clear what is happening from the dump.
405406
OS << " clause";

‎clang/lib/Parse/ParseOpenACC.cpp

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -960,7 +960,8 @@ Parser::OpenACCClauseParseResult Parser::ParseOpenACCClauseParams(
960960

961961
// TODO OpenACC: as we implement the 'rest' of the above, this 'if' should
962962
// be removed leaving just the 'setIntExprDetails'.
963-
if (ClauseKind == OpenACCClauseKind::NumWorkers)
963+
if (ClauseKind == OpenACCClauseKind::NumWorkers ||
964+
ClauseKind == OpenACCClauseKind::VectorLength)
964965
ParsedClause.setIntExprDetails(IntExpr.get());
965966

966967
break;

‎clang/lib/Sema/SemaOpenACC.cpp

Lines changed: 20 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -92,6 +92,7 @@ bool doesClauseApplyToDirective(OpenACCDirectiveKind DirectiveKind,
9292
return false;
9393
}
9494
case OpenACCClauseKind::NumWorkers:
95+
case OpenACCClauseKind::VectorLength:
9596
switch (DirectiveKind) {
9697
case OpenACCDirectiveKind::Parallel:
9798
case OpenACCDirectiveKind::Kernels:
@@ -248,6 +249,25 @@ SemaOpenACC::ActOnClause(ArrayRef<const OpenACCClause *> ExistingClauses,
248249
getASTContext(), Clause.getBeginLoc(), Clause.getLParenLoc(),
249250
Clause.getIntExprs()[0], Clause.getEndLoc());
250251
}
252+
case OpenACCClauseKind::VectorLength: {
253+
// Restrictions only properly implemented on 'compute' constructs, and
254+
// 'compute' constructs are the only construct that can do anything with
255+
// this yet, so skip/treat as unimplemented in this case.
256+
if (!isOpenACCComputeDirectiveKind(Clause.getDirectiveKind()))
257+
break;
258+
259+
// There is no prose in the standard that says duplicates aren't allowed,
260+
// but this diagnostic is present in other compilers, as well as makes
261+
// sense.
262+
if (checkAlreadyHasClauseOfKind(*this, ExistingClauses, Clause))
263+
return nullptr;
264+
265+
assert(Clause.getIntExprs().size() == 1 &&
266+
"Invalid number of expressions for VectorLength");
267+
return OpenACCVectorLengthClause::Create(
268+
getASTContext(), Clause.getBeginLoc(), Clause.getLParenLoc(),
269+
Clause.getIntExprs()[0], Clause.getEndLoc());
270+
}
251271
default:
252272
break;
253273
}

‎clang/lib/Sema/TreeTransform.h

Lines changed: 23 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -11181,6 +11181,29 @@ void OpenACCClauseTransform<Derived>::VisitNumWorkersClause(
1118111181
ParsedClause.getLParenLoc(), ParsedClause.getIntExprs()[0],
1118211182
ParsedClause.getEndLoc());
1118311183
}
11184+
11185+
template <typename Derived>
11186+
void OpenACCClauseTransform<Derived>::VisitVectorLengthClause(
11187+
const OpenACCVectorLengthClause &C) {
11188+
Expr *IntExpr = const_cast<Expr *>(C.getIntExpr());
11189+
assert(IntExpr && "vector_length clause constructed with invalid int expr");
11190+
11191+
ExprResult Res = Self.TransformExpr(IntExpr);
11192+
if (!Res.isUsable())
11193+
return;
11194+
11195+
Res = Self.getSema().OpenACC().ActOnIntExpr(OpenACCDirectiveKind::Invalid,
11196+
C.getClauseKind(),
11197+
C.getBeginLoc(), Res.get());
11198+
if (!Res.isUsable())
11199+
return;
11200+
11201+
ParsedClause.setIntExprDetails(Res.get());
11202+
NewClause = OpenACCVectorLengthClause::Create(
11203+
Self.getSema().getASTContext(), ParsedClause.getBeginLoc(),
11204+
ParsedClause.getLParenLoc(), ParsedClause.getIntExprs()[0],
11205+
ParsedClause.getEndLoc());
11206+
}
1118411207
} // namespace
1118511208
template <typename Derived>
1118611209
OpenACCClause *TreeTransform<Derived>::TransformOpenACCClause(

‎clang/lib/Serialization/ASTReader.cpp

Lines changed: 6 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -11792,6 +11792,12 @@ OpenACCClause *ASTRecordReader::readOpenACCClause() {
1179211792
return OpenACCNumWorkersClause::Create(getContext(), BeginLoc, LParenLoc,
1179311793
IntExpr, EndLoc);
1179411794
}
11795+
case OpenACCClauseKind::VectorLength: {
11796+
SourceLocation LParenLoc = readSourceLocation();
11797+
Expr *IntExpr = readSubExpr();
11798+
return OpenACCVectorLengthClause::Create(getContext(), BeginLoc, LParenLoc,
11799+
IntExpr, EndLoc);
11800+
}
1179511801
case OpenACCClauseKind::Finalize:
1179611802
case OpenACCClauseKind::IfPresent:
1179711803
case OpenACCClauseKind::Seq:
@@ -11820,7 +11826,6 @@ OpenACCClause *ASTRecordReader::readOpenACCClause() {
1182011826
case OpenACCClauseKind::Reduction:
1182111827
case OpenACCClauseKind::Collapse:
1182211828
case OpenACCClauseKind::Bind:
11823-
case OpenACCClauseKind::VectorLength:
1182411829
case OpenACCClauseKind::NumGangs:
1182511830
case OpenACCClauseKind::DeviceNum:
1182611831
case OpenACCClauseKind::DefaultAsync:

‎clang/lib/Serialization/ASTWriter.cpp

Lines changed: 8 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -5014,7 +5014,7 @@ void ASTWriter::WriteSpecialDeclRecords(Sema &SemaRef) {
50145014

50155015
if (!ModularCodegenDecls.empty())
50165016
Stream.EmitRecord(MODULAR_CODEGEN_DECLS, ModularCodegenDecls);
5017-
5017+
50185018
// Write the record containing tentative definitions.
50195019
RecordData TentativeDefinitions;
50205020
AddLazyVectorEmiitedDecls(*this, SemaRef.TentativeDefinitions,
@@ -5135,7 +5135,7 @@ void ASTWriter::WriteSpecialDeclRecords(Sema &SemaRef) {
51355135
}
51365136
if (!UndefinedButUsed.empty())
51375137
Stream.EmitRecord(UNDEFINED_BUT_USED, UndefinedButUsed);
5138-
5138+
51395139
// Write all delete-expressions that we would like to
51405140
// analyze later in AST.
51415141
RecordData DeleteExprsToAnalyze;
@@ -7663,6 +7663,12 @@ void ASTRecordWriter::writeOpenACCClause(const OpenACCClause *C) {
76637663
AddStmt(const_cast<Expr *>(NWC->getIntExpr()));
76647664
return;
76657665
}
7666+
case OpenACCClauseKind::VectorLength: {
7667+
const auto *NWC = cast<OpenACCVectorLengthClause>(C);
7668+
writeSourceLocation(NWC->getLParenLoc());
7669+
AddStmt(const_cast<Expr *>(NWC->getIntExpr()));
7670+
return;
7671+
}
76667672
case OpenACCClauseKind::Finalize:
76677673
case OpenACCClauseKind::IfPresent:
76687674
case OpenACCClauseKind::Seq:
@@ -7691,7 +7697,6 @@ void ASTRecordWriter::writeOpenACCClause(const OpenACCClause *C) {
76917697
case OpenACCClauseKind::Reduction:
76927698
case OpenACCClauseKind::Collapse:
76937699
case OpenACCClauseKind::Bind:
7694-
case OpenACCClauseKind::VectorLength:
76957700
case OpenACCClauseKind::NumGangs:
76967701
case OpenACCClauseKind::DeviceNum:
76977702
case OpenACCClauseKind::DefaultAsync:

‎clang/test/ParserOpenACC/parse-clauses.c

Lines changed: 0 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -893,11 +893,9 @@ void IntExprParsing() {
893893
#pragma acc parallel vector_length(5, 4)
894894
{}
895895

896-
// expected-warning@+1{{OpenACC clause 'vector_length' not yet implemented, clause ignored}}
897896
#pragma acc parallel vector_length(5)
898897
{}
899898

900-
// expected-warning@+1{{OpenACC clause 'vector_length' not yet implemented, clause ignored}}
901899
#pragma acc parallel vector_length(returns_int())
902900
{}
903901

‎clang/test/ParserOpenACC/parse-clauses.cpp

Lines changed: 0 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -12,11 +12,9 @@ void templ() {
1212
#pragma acc loop collapse(T::value)
1313
for(;;){}
1414

15-
// expected-warning@+1{{OpenACC clause 'vector_length' not yet implemented, clause ignored}}
1615
#pragma acc parallel vector_length(T::value)
1716
for(;;){}
1817

19-
// expected-warning@+1{{OpenACC clause 'vector_length' not yet implemented, clause ignored}}
2018
#pragma acc parallel vector_length(I)
2119
for(;;){}
2220

‎clang/test/SemaOpenACC/compute-construct-intexpr-clause-ast.cpp

Lines changed: 49 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -77,6 +77,17 @@ void NormalUses() {
7777
// CHECK-NEXT: WhileStmt
7878
// CHECK-NEXT: CXXBoolLiteralExpr
7979
// CHECK-NEXT: CompoundStmt
80+
81+
#pragma acc kernels vector_length(some_short())
82+
while(true){}
83+
// CHECK-NEXT: OpenACCComputeConstruct{{.*}}kernels
84+
// CHECK-NEXT: vector_length clause
85+
// CHECK-NEXT: CallExpr{{.*}}'short'
86+
// CHECK-NEXT: ImplicitCastExpr{{.*}}'short (*)()' <FunctionToPointerDecay>
87+
// CHECK-NEXT: DeclRefExpr{{.*}}'short ()' lvalue Function{{.*}} 'some_short' 'short ()'
88+
// CHECK-NEXT: WhileStmt
89+
// CHECK-NEXT: CXXBoolLiteralExpr
90+
// CHECK-NEXT: CompoundStmt
8091
}
8192

8293
template<typename T, typename U>
@@ -157,6 +168,25 @@ void TemplUses(T t, U u) {
157168
// CHECK-NEXT: CXXBoolLiteralExpr
158169
// CHECK-NEXT: CompoundStmt
159170

171+
#pragma acc kernels vector_length(u)
172+
while(true){}
173+
// CHECK-NEXT: OpenACCComputeConstruct{{.*}}kernels
174+
// CHECK-NEXT: vector_length clause
175+
// CHECK-NEXT: DeclRefExpr{{.*}} 'U' lvalue ParmVar{{.*}} 'u' 'U'
176+
// CHECK-NEXT: WhileStmt
177+
// CHECK-NEXT: CXXBoolLiteralExpr
178+
// CHECK-NEXT: CompoundStmt
179+
180+
#pragma acc parallel vector_length(U::value)
181+
while(true){}
182+
// CHECK-NEXT: OpenACCComputeConstruct{{.*}}parallel
183+
// CHECK-NEXT: vector_length clause
184+
// CHECK-NEXT: DependentScopeDeclRefExpr{{.*}} '<dependent type>' lvalue
185+
// CHECK-NEXT: NestedNameSpecifier TypeSpec 'U'
186+
// CHECK-NEXT: WhileStmt
187+
// CHECK-NEXT: CXXBoolLiteralExpr
188+
// CHECK-NEXT: CompoundStmt
189+
160190
// Check the instantiated versions of the above.
161191
// CHECK-NEXT: FunctionDecl{{.*}} used TemplUses 'void (CorrectConvert, HasInt)' implicit_instantiation
162192
// CHECK-NEXT: TemplateArgument type 'CorrectConvert'
@@ -239,6 +269,25 @@ void TemplUses(T t, U u) {
239269
// CHECK-NEXT: WhileStmt
240270
// CHECK-NEXT: CXXBoolLiteralExpr
241271
// CHECK-NEXT: CompoundStmt
272+
273+
// CHECK-NEXT: OpenACCComputeConstruct{{.*}}kernels
274+
// CHECK-NEXT: vector_length clause
275+
// CHECK-NEXT: ImplicitCastExpr{{.*}} 'char' <UserDefinedConversion>
276+
// CHECK-NEXT: CXXMemberCallExpr{{.*}}'char'
277+
// CHECK-NEXT: MemberExpr{{.*}} '<bound member function type>' .operator char
278+
// CHECK-NEXT: DeclRefExpr{{.*}} 'HasInt' lvalue ParmVar
279+
// CHECK-NEXT: WhileStmt
280+
// CHECK-NEXT: CXXBoolLiteralExpr
281+
// CHECK-NEXT: CompoundStmt
282+
283+
// CHECK-NEXT: OpenACCComputeConstruct{{.*}}parallel
284+
// CHECK-NEXT: vector_length clause
285+
// CHECK-NEXT: ImplicitCastExpr{{.*}} 'int' <LValueToRValue>
286+
// CHECK-NEXT: DeclRefExpr{{.*}} 'const int' lvalue Var{{.*}} 'value' 'const int'
287+
// CHECK-NEXT: NestedNameSpecifier TypeSpec 'HasInt'
288+
// CHECK-NEXT: WhileStmt
289+
// CHECK-NEXT: CXXBoolLiteralExpr
290+
// CHECK-NEXT: CompoundStmt
242291
}
243292

244293
struct HasInt {

‎clang/test/SemaOpenACC/compute-construct-vector_length-clause.c

Lines changed: 33 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,33 @@
1+
// RUN: %clang_cc1 %s -fopenacc -verify
2+
3+
short getS();
4+
5+
void Test() {
6+
#pragma acc parallel vector_length(1)
7+
while(1);
8+
#pragma acc kernels vector_length(1)
9+
while(1);
10+
11+
// expected-error@+1{{OpenACC 'vector_length' clause is not valid on 'serial' directive}}
12+
#pragma acc serial vector_length(1)
13+
while(1);
14+
15+
struct NotConvertible{} NC;
16+
// expected-error@+1{{OpenACC clause 'vector_length' requires expression of integer type ('struct NotConvertible' invalid)}}
17+
#pragma acc parallel vector_length(NC)
18+
while(1);
19+
20+
#pragma acc kernels vector_length(getS())
21+
while(1);
22+
23+
struct Incomplete *SomeIncomplete;
24+
25+
// expected-error@+1{{OpenACC clause 'vector_length' requires expression of integer type ('struct Incomplete' invalid)}}
26+
#pragma acc kernels vector_length(*SomeIncomplete)
27+
while(1);
28+
29+
enum E{A} SomeE;
30+
31+
#pragma acc kernels vector_length(SomeE)
32+
while(1);
33+
}

0 commit comments

Comments
 (0)