1 //===--- ParseOpenACC.cpp - OpenACC-specific parsing support --------------===//
2 //
3 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4 // See https://llvm.org/LICENSE.txt for license information.
5 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6 //
7 //===----------------------------------------------------------------------===//
8 //
9 // This file implements the parsing logic for OpenACC language features.
10 //
11 //===----------------------------------------------------------------------===//
12 
13 #include "clang/Basic/OpenACCKinds.h"
14 #include "clang/Parse/ParseDiagnostic.h"
15 #include "clang/Parse/Parser.h"
16 #include "clang/Parse/RAIIObjectsForParser.h"
17 #include "llvm/ADT/StringRef.h"
18 #include "llvm/ADT/StringSwitch.h"
19 
20 using namespace clang;
21 using namespace llvm;
22 
23 namespace {
24 // An enum that contains the extended 'partial' parsed variants. This type
25 // should never escape the initial parse functionality, but is useful for
26 // simplifying the implementation.
27 enum class OpenACCDirectiveKindEx {
28   Invalid = static_cast<int>(OpenACCDirectiveKind::Invalid),
29   // 'enter data' and 'exit data'
30   Enter,
31   Exit,
32 };
33 
34 // Translate single-token string representations to the OpenACC Directive Kind.
35 // This doesn't completely comprehend 'Compound Constructs' (as it just
36 // identifies the first token), and doesn't fully handle 'enter data', 'exit
37 // data', nor any of the 'atomic' variants, just the first token of each.  So
38 // this should only be used by `ParseOpenACCDirectiveKind`.
39 OpenACCDirectiveKindEx getOpenACCDirectiveKind(Token Tok) {
40   if (!Tok.is(tok::identifier))
41     return OpenACCDirectiveKindEx::Invalid;
42   OpenACCDirectiveKind DirKind =
43       llvm::StringSwitch<OpenACCDirectiveKind>(
44           Tok.getIdentifierInfo()->getName())
45           .Case("parallel", OpenACCDirectiveKind::Parallel)
46           .Case("serial", OpenACCDirectiveKind::Serial)
47           .Case("kernels", OpenACCDirectiveKind::Kernels)
48           .Case("data", OpenACCDirectiveKind::Data)
49           .Case("host_data", OpenACCDirectiveKind::HostData)
50           .Case("loop", OpenACCDirectiveKind::Loop)
51           .Case("cache", OpenACCDirectiveKind::Cache)
52           .Case("atomic", OpenACCDirectiveKind::Atomic)
53           .Case("routine", OpenACCDirectiveKind::Routine)
54           .Case("declare", OpenACCDirectiveKind::Declare)
55           .Case("init", OpenACCDirectiveKind::Init)
56           .Case("shutdown", OpenACCDirectiveKind::Shutdown)
57           .Case("set", OpenACCDirectiveKind::Shutdown)
58           .Case("update", OpenACCDirectiveKind::Update)
59           .Case("wait", OpenACCDirectiveKind::Wait)
60           .Default(OpenACCDirectiveKind::Invalid);
61 
62   if (DirKind != OpenACCDirectiveKind::Invalid)
63     return static_cast<OpenACCDirectiveKindEx>(DirKind);
64 
65   return llvm::StringSwitch<OpenACCDirectiveKindEx>(
66              Tok.getIdentifierInfo()->getName())
67       .Case("enter", OpenACCDirectiveKindEx::Enter)
68       .Case("exit", OpenACCDirectiveKindEx::Exit)
69       .Default(OpenACCDirectiveKindEx::Invalid);
70 }
71 
72 // Translate single-token string representations to the OpenCC Clause Kind.
73 OpenACCClauseKind getOpenACCClauseKind(Token Tok) {
74   // auto is a keyword in some language modes, so make sure we parse it
75   // correctly.
76   if (Tok.is(tok::kw_auto))
77     return OpenACCClauseKind::Auto;
78 
79   if (!Tok.is(tok::identifier))
80     return OpenACCClauseKind::Invalid;
81 
82   return llvm::StringSwitch<OpenACCClauseKind>(
83              Tok.getIdentifierInfo()->getName())
84       .Case("auto", OpenACCClauseKind::Auto)
85       .Case("finalize", OpenACCClauseKind::Finalize)
86       .Case("if_present", OpenACCClauseKind::IfPresent)
87       .Case("independent", OpenACCClauseKind::Independent)
88       .Case("nohost", OpenACCClauseKind::NoHost)
89       .Case("seq", OpenACCClauseKind::Seq)
90       .Case("vector", OpenACCClauseKind::Vector)
91       .Case("worker", OpenACCClauseKind::Worker)
92       .Default(OpenACCClauseKind::Invalid);
93 }
94 
95 // Since 'atomic' is effectively a compound directive, this will decode the
96 // second part of the directive.
97 OpenACCAtomicKind getOpenACCAtomicKind(Token Tok) {
98   if (!Tok.is(tok::identifier))
99     return OpenACCAtomicKind::Invalid;
100   return llvm::StringSwitch<OpenACCAtomicKind>(
101              Tok.getIdentifierInfo()->getName())
102       .Case("read", OpenACCAtomicKind::Read)
103       .Case("write", OpenACCAtomicKind::Write)
104       .Case("update", OpenACCAtomicKind::Update)
105       .Case("capture", OpenACCAtomicKind::Capture)
106       .Default(OpenACCAtomicKind::Invalid);
107 }
108 
109 enum class OpenACCSpecialTokenKind {
110   ReadOnly,
111   DevNum,
112   Queues,
113 };
114 
115 bool isOpenACCSpecialToken(OpenACCSpecialTokenKind Kind, Token Tok) {
116   if (!Tok.is(tok::identifier))
117     return false;
118 
119   switch (Kind) {
120   case OpenACCSpecialTokenKind::ReadOnly:
121     return Tok.getIdentifierInfo()->isStr("readonly");
122   case OpenACCSpecialTokenKind::DevNum:
123     return Tok.getIdentifierInfo()->isStr("devnum");
124   case OpenACCSpecialTokenKind::Queues:
125     return Tok.getIdentifierInfo()->isStr("queues");
126   }
127   llvm_unreachable("Unknown 'Kind' Passed");
128 }
129 
130 bool isOpenACCDirectiveKind(OpenACCDirectiveKind Kind, Token Tok) {
131   if (!Tok.is(tok::identifier))
132     return false;
133 
134   switch (Kind) {
135   case OpenACCDirectiveKind::Parallel:
136     return Tok.getIdentifierInfo()->isStr("parallel");
137   case OpenACCDirectiveKind::Serial:
138     return Tok.getIdentifierInfo()->isStr("serial");
139   case OpenACCDirectiveKind::Kernels:
140     return Tok.getIdentifierInfo()->isStr("kernels");
141   case OpenACCDirectiveKind::Data:
142     return Tok.getIdentifierInfo()->isStr("data");
143   case OpenACCDirectiveKind::HostData:
144     return Tok.getIdentifierInfo()->isStr("host_data");
145   case OpenACCDirectiveKind::Loop:
146     return Tok.getIdentifierInfo()->isStr("loop");
147   case OpenACCDirectiveKind::Cache:
148     return Tok.getIdentifierInfo()->isStr("cache");
149 
150   case OpenACCDirectiveKind::ParallelLoop:
151   case OpenACCDirectiveKind::SerialLoop:
152   case OpenACCDirectiveKind::KernelsLoop:
153   case OpenACCDirectiveKind::EnterData:
154   case OpenACCDirectiveKind::ExitData:
155     return false;
156 
157   case OpenACCDirectiveKind::Atomic:
158     return Tok.getIdentifierInfo()->isStr("atomic");
159   case OpenACCDirectiveKind::Routine:
160     return Tok.getIdentifierInfo()->isStr("routine");
161   case OpenACCDirectiveKind::Declare:
162     return Tok.getIdentifierInfo()->isStr("declare");
163   case OpenACCDirectiveKind::Init:
164     return Tok.getIdentifierInfo()->isStr("init");
165   case OpenACCDirectiveKind::Shutdown:
166     return Tok.getIdentifierInfo()->isStr("shutdown");
167   case OpenACCDirectiveKind::Set:
168     return Tok.getIdentifierInfo()->isStr("set");
169   case OpenACCDirectiveKind::Update:
170     return Tok.getIdentifierInfo()->isStr("update");
171   case OpenACCDirectiveKind::Wait:
172     return Tok.getIdentifierInfo()->isStr("wait");
173   case OpenACCDirectiveKind::Invalid:
174     return false;
175   }
176   llvm_unreachable("Unknown 'Kind' Passed");
177 }
178 
179 OpenACCDirectiveKind
180 ParseOpenACCEnterExitDataDirective(Parser &P, Token FirstTok,
181                                    OpenACCDirectiveKindEx ExtDirKind) {
182   Token SecondTok = P.getCurToken();
183 
184   if (SecondTok.isAnnotation()) {
185     P.Diag(FirstTok, diag::err_acc_invalid_directive)
186         << 0 << FirstTok.getIdentifierInfo();
187     return OpenACCDirectiveKind::Invalid;
188   }
189 
190   // Consume the second name anyway, this way we can continue on without making
191   // this oddly look like a clause.
192   P.ConsumeAnyToken();
193 
194   if (!isOpenACCDirectiveKind(OpenACCDirectiveKind::Data, SecondTok)) {
195     if (!SecondTok.is(tok::identifier))
196       P.Diag(SecondTok, diag::err_expected) << tok::identifier;
197     else
198       P.Diag(FirstTok, diag::err_acc_invalid_directive)
199           << 1 << FirstTok.getIdentifierInfo()->getName()
200           << SecondTok.getIdentifierInfo()->getName();
201     return OpenACCDirectiveKind::Invalid;
202   }
203 
204   return ExtDirKind == OpenACCDirectiveKindEx::Enter
205              ? OpenACCDirectiveKind::EnterData
206              : OpenACCDirectiveKind::ExitData;
207 }
208 
209 OpenACCAtomicKind ParseOpenACCAtomicKind(Parser &P) {
210   Token AtomicClauseToken = P.getCurToken();
211 
212   // #pragma acc atomic is equivilent to update:
213   if (AtomicClauseToken.isAnnotation())
214     return OpenACCAtomicKind::Update;
215 
216   OpenACCAtomicKind AtomicKind = getOpenACCAtomicKind(AtomicClauseToken);
217 
218   // If we don't know what this is, treat it as 'nothing', and treat the rest of
219   // this as a clause list, which, despite being invalid, is likely what the
220   // user was trying to do.
221   if (AtomicKind == OpenACCAtomicKind::Invalid)
222     return OpenACCAtomicKind::Update;
223 
224   P.ConsumeToken();
225   return AtomicKind;
226 }
227 
228 // Parse and consume the tokens for OpenACC Directive/Construct kinds.
229 OpenACCDirectiveKind ParseOpenACCDirectiveKind(Parser &P) {
230   Token FirstTok = P.getCurToken();
231 
232   // Just #pragma acc can get us immediately to the end, make sure we don't
233   // introspect on the spelling before then.
234   if (FirstTok.isNot(tok::identifier)) {
235     P.Diag(FirstTok, diag::err_acc_missing_directive);
236 
237     if (P.getCurToken().isNot(tok::annot_pragma_openacc_end))
238       P.ConsumeAnyToken();
239 
240     return OpenACCDirectiveKind::Invalid;
241   }
242 
243   P.ConsumeToken();
244 
245   OpenACCDirectiveKindEx ExDirKind = getOpenACCDirectiveKind(FirstTok);
246 
247   // OpenACCDirectiveKindEx is meant to be an extended list
248   // over OpenACCDirectiveKind, so any value below Invalid is one of the
249   // OpenACCDirectiveKind values.  This switch takes care of all of the extra
250   // parsing required for the Extended values.  At the end of this block,
251   // ExDirKind can be assumed to be a valid OpenACCDirectiveKind, so we can
252   // immediately cast it and use it as that.
253   if (ExDirKind >= OpenACCDirectiveKindEx::Invalid) {
254     switch (ExDirKind) {
255     case OpenACCDirectiveKindEx::Invalid: {
256       P.Diag(FirstTok, diag::err_acc_invalid_directive)
257           << 0 << FirstTok.getIdentifierInfo();
258       return OpenACCDirectiveKind::Invalid;
259     }
260     case OpenACCDirectiveKindEx::Enter:
261     case OpenACCDirectiveKindEx::Exit:
262       return ParseOpenACCEnterExitDataDirective(P, FirstTok, ExDirKind);
263     }
264   }
265 
266   OpenACCDirectiveKind DirKind = static_cast<OpenACCDirectiveKind>(ExDirKind);
267 
268   // Combined Constructs allows parallel loop, serial loop, or kernels loop. Any
269   // other attempt at a combined construct will be diagnosed as an invalid
270   // clause.
271   Token SecondTok = P.getCurToken();
272   if (!SecondTok.isAnnotation() &&
273       isOpenACCDirectiveKind(OpenACCDirectiveKind::Loop, SecondTok)) {
274     switch (DirKind) {
275     default:
276       // Nothing to do except in the below cases, as they should be diagnosed as
277       // a clause.
278       break;
279     case OpenACCDirectiveKind::Parallel:
280       P.ConsumeToken();
281       return OpenACCDirectiveKind::ParallelLoop;
282     case OpenACCDirectiveKind::Serial:
283       P.ConsumeToken();
284       return OpenACCDirectiveKind::SerialLoop;
285     case OpenACCDirectiveKind::Kernels:
286       P.ConsumeToken();
287       return OpenACCDirectiveKind::KernelsLoop;
288     }
289   }
290 
291   return DirKind;
292 }
293 
294 // The OpenACC Clause List is a comma or space-delimited list of clauses (see
295 // the comment on ParseOpenACCClauseList).  The concept of a 'clause' doesn't
296 // really have its owner grammar and each individual one has its own definition.
297 // However, they all are named with a single-identifier (or auto!) token,
298 // followed in some cases by either braces or parens.
299 bool ParseOpenACCClause(Parser &P) {
300   if (!P.getCurToken().isOneOf(tok::identifier, tok::kw_auto))
301     return P.Diag(P.getCurToken(), diag::err_expected) << tok::identifier;
302 
303   OpenACCClauseKind Kind = getOpenACCClauseKind(P.getCurToken());
304 
305   if (Kind == OpenACCClauseKind::Invalid)
306     return P.Diag(P.getCurToken(), diag::err_acc_invalid_clause)
307            << P.getCurToken().getIdentifierInfo();
308 
309   // Consume the clause name.
310   P.ConsumeToken();
311 
312   // FIXME: For future clauses, we need to handle parens/etc below.
313   return false;
314 }
315 
316 // Skip until we see the end of pragma token, but don't consume it. This is us
317 // just giving up on the rest of the pragma so we can continue executing. We
318 // have to do this because 'SkipUntil' considers paren balancing, which isn't
319 // what we want.
320 void SkipUntilEndOfDirective(Parser &P) {
321   while (P.getCurToken().isNot(tok::annot_pragma_openacc_end))
322     P.ConsumeAnyToken();
323 }
324 
325 // OpenACC 3.3, section 1.7:
326 // To simplify the specification and convey appropriate constraint information,
327 // a pqr-list is a comma-separated list of pdr items. The one exception is a
328 // clause-list, which is a list of one or more clauses optionally separated by
329 // commas.
330 void ParseOpenACCClauseList(Parser &P) {
331   bool FirstClause = true;
332   while (P.getCurToken().isNot(tok::annot_pragma_openacc_end)) {
333     // Comma is optional in a clause-list.
334     if (!FirstClause && P.getCurToken().is(tok::comma))
335       P.ConsumeToken();
336     FirstClause = false;
337 
338     // Recovering from a bad clause is really difficult, so we just give up on
339     // error.
340     if (ParseOpenACCClause(P)) {
341       SkipUntilEndOfDirective(P);
342       return;
343     }
344   }
345 }
346 
347 } // namespace
348 
349 /// OpenACC 3.3, section 2.16:
350 /// In this section and throughout the specification, the term wait-argument
351 /// means:
352 /// [ devnum : int-expr : ] [ queues : ] async-argument-list
353 bool Parser::ParseOpenACCWaitArgument() {
354   // [devnum : int-expr : ]
355   if (isOpenACCSpecialToken(OpenACCSpecialTokenKind::DevNum, Tok) &&
356       NextToken().is(tok::colon)) {
357     // Consume devnum.
358     ConsumeToken();
359     // Consume colon.
360     ConsumeToken();
361 
362     ExprResult IntExpr =
363         getActions().CorrectDelayedTyposInExpr(ParseAssignmentExpression());
364     if (IntExpr.isInvalid())
365       return true;
366 
367     if (ExpectAndConsume(tok::colon))
368       return true;
369   }
370 
371   // [ queues : ]
372   if (isOpenACCSpecialToken(OpenACCSpecialTokenKind::Queues, Tok) &&
373       NextToken().is(tok::colon)) {
374     // Consume queues.
375     ConsumeToken();
376     // Consume colon.
377     ConsumeToken();
378   }
379 
380   // OpenACC 3.3, section 2.16:
381   // the term 'async-argument' means a nonnegative scalar integer expression, or
382   // one of the special values 'acc_async_noval' or 'acc_async_sync', as defined
383   // in the C header file and the Fortran opacc module.
384   //
385   // We are parsing this simply as list of assignment expressions (to avoid
386   // comma being troublesome), and will ensure it is an integral type.  The
387   // 'special' types are defined as macros, so we can't really check those
388   // (other than perhaps as values at one point?), but the standard does say it
389   // is implementation-defined to use any other negative value.
390   //
391   //
392   bool FirstArg = true;
393   while (!getCurToken().isOneOf(tok::r_paren, tok::annot_pragma_openacc_end)) {
394     if (!FirstArg) {
395       if (ExpectAndConsume(tok::comma))
396         return true;
397     }
398     FirstArg = false;
399 
400     ExprResult CurArg =
401         getActions().CorrectDelayedTyposInExpr(ParseAssignmentExpression());
402 
403     if (CurArg.isInvalid())
404       return true;
405   }
406 
407   return false;
408 }
409 
410 ExprResult Parser::ParseOpenACCIDExpression() {
411   ExprResult Res;
412   if (getLangOpts().CPlusPlus) {
413     Res = ParseCXXIdExpression(/*isAddressOfOperand=*/false);
414   } else {
415     // There isn't anything quite the same as ParseCXXIdExpression for C, so we
416     // need to get the identifier, then call into Sema ourselves.
417 
418     if (Tok.isNot(tok::identifier)) {
419       Diag(Tok, diag::err_expected) << tok::identifier;
420       return ExprError();
421     }
422 
423     Token FuncName = getCurToken();
424     UnqualifiedId Name;
425     CXXScopeSpec ScopeSpec;
426     SourceLocation TemplateKWLoc;
427     Name.setIdentifier(FuncName.getIdentifierInfo(), ConsumeToken());
428 
429     // Ensure this is a valid identifier. We don't accept causing implicit
430     // function declarations per the spec, so always claim to not have trailing
431     // L Paren.
432     Res = Actions.ActOnIdExpression(getCurScope(), ScopeSpec, TemplateKWLoc,
433                                     Name, /*HasTrailingLParen=*/false,
434                                     /*isAddressOfOperand=*/false);
435   }
436 
437   return getActions().CorrectDelayedTyposInExpr(Res);
438 }
439 
440 /// OpenACC 3.3, section 2.10:
441 /// A 'var' in a cache directive must be a single array element or a simple
442 /// subarray.  In C and C++, a simple subarray is an array name followed by an
443 /// extended array range specification in brackets, with a start and length such
444 /// as:
445 ///
446 /// arr[lower:length]
447 ///
448 bool Parser::ParseOpenACCCacheVar() {
449   ExprResult ArrayName = ParseOpenACCIDExpression();
450   if (ArrayName.isInvalid())
451     return true;
452 
453   // If the expression is invalid, just continue parsing the brackets, there
454   // is likely other useful diagnostics we can emit inside of those.
455 
456   BalancedDelimiterTracker SquareBrackets(*this, tok::l_square,
457                                           tok::annot_pragma_openacc_end);
458 
459   // Square brackets are required, so error here, and try to recover by moving
460   // until the next comma, or the close paren/end of pragma.
461   if (SquareBrackets.expectAndConsume()) {
462     SkipUntil(tok::comma, tok::r_paren, tok::annot_pragma_openacc_end,
463               Parser::StopBeforeMatch);
464     return true;
465   }
466 
467   ExprResult Lower = getActions().CorrectDelayedTyposInExpr(ParseExpression());
468   if (Lower.isInvalid())
469     return true;
470 
471   // The 'length' expression is optional, as this could be a single array
472   // element. If there is no colon, we can treat it as that.
473   if (getCurToken().is(tok::colon)) {
474     ConsumeToken();
475     ExprResult Length =
476         getActions().CorrectDelayedTyposInExpr(ParseExpression());
477     if (Length.isInvalid())
478       return true;
479   }
480 
481   // Diagnose the square bracket being in the wrong place and continue.
482   return SquareBrackets.consumeClose();
483 }
484 
485 /// OpenACC 3.3, section 2.10:
486 /// In C and C++, the syntax of the cache directive is:
487 ///
488 /// #pragma acc cache ([readonly:]var-list) new-line
489 void Parser::ParseOpenACCCacheVarList() {
490   // If this is the end of the line, just return 'false' and count on the close
491   // paren diagnostic to catch the issue.
492   if (getCurToken().isAnnotation())
493     return;
494 
495   // The VarList is an optional `readonly:` followed by a list of a variable
496   // specifications.  First, see if we have `readonly:`, else we back-out and
497   // treat it like the beginning of a reference to a potentially-existing
498   // `readonly` variable.
499   if (isOpenACCSpecialToken(OpenACCSpecialTokenKind::ReadOnly, Tok) &&
500       NextToken().is(tok::colon)) {
501     // Consume both tokens.
502     ConsumeToken();
503     ConsumeToken();
504     // FIXME: Record that this is a 'readonly' so that we can use that during
505     // Sema/AST generation.
506   }
507 
508   bool FirstArray = true;
509   while (!getCurToken().isOneOf(tok::r_paren, tok::annot_pragma_openacc_end)) {
510     if (!FirstArray)
511       ExpectAndConsume(tok::comma);
512     FirstArray = false;
513     if (ParseOpenACCCacheVar())
514       SkipUntil(tok::r_paren, tok::annot_pragma_openacc_end, tok::comma,
515                 StopBeforeMatch);
516   }
517 }
518 
519 void Parser::ParseOpenACCDirective() {
520   OpenACCDirectiveKind DirKind = ParseOpenACCDirectiveKind(*this);
521 
522   // Once we've parsed the construct/directive name, some have additional
523   // specifiers that need to be taken care of. Atomic has an 'atomic-clause'
524   // that needs to be parsed.
525   if (DirKind == OpenACCDirectiveKind::Atomic)
526     ParseOpenACCAtomicKind(*this);
527 
528   // We've successfully parsed the construct/directive name, however a few of
529   // the constructs have optional parens that contain further details.
530   BalancedDelimiterTracker T(*this, tok::l_paren,
531                              tok::annot_pragma_openacc_end);
532 
533   if (!T.consumeOpen()) {
534     switch (DirKind) {
535     default:
536       Diag(T.getOpenLocation(), diag::err_acc_invalid_open_paren);
537       T.skipToEnd();
538       break;
539     case OpenACCDirectiveKind::Routine: {
540       // Routine has an optional paren-wrapped name of a function in the local
541       // scope. We parse the name, emitting any diagnostics
542       ExprResult RoutineName = ParseOpenACCIDExpression();
543       // If the routine name is invalid, just skip until the closing paren to
544       // recover more gracefully.
545       if (RoutineName.isInvalid())
546         T.skipToEnd();
547       else
548         T.consumeClose();
549       break;
550     }
551     case OpenACCDirectiveKind::Cache:
552       ParseOpenACCCacheVarList();
553       // The ParseOpenACCCacheVarList function manages to recover from failures,
554       // so we can always consume the close.
555       T.consumeClose();
556       break;
557     case OpenACCDirectiveKind::Wait:
558       // OpenACC has an optional paren-wrapped 'wait-argument'.
559       if (ParseOpenACCWaitArgument())
560         T.skipToEnd();
561       else
562         T.consumeClose();
563       break;
564     }
565   } else if (DirKind == OpenACCDirectiveKind::Cache) {
566     // Cache's paren var-list is required, so error here if it isn't provided.
567     // We know that the consumeOpen above left the first non-paren here, so
568     // diagnose, then continue as if it was completely omitted.
569     Diag(Tok, diag::err_expected) << tok::l_paren;
570   }
571 
572   // Parses the list of clauses, if present.
573   ParseOpenACCClauseList(*this);
574 
575   Diag(getCurToken(), diag::warn_pragma_acc_unimplemented);
576   assert(Tok.is(tok::annot_pragma_openacc_end) &&
577          "Didn't parse all OpenACC Clauses");
578   ConsumeAnnotationToken();
579 }
580 
581 // Parse OpenACC directive on a declaration.
582 Parser::DeclGroupPtrTy Parser::ParseOpenACCDirectiveDecl() {
583   assert(Tok.is(tok::annot_pragma_openacc) && "expected OpenACC Start Token");
584 
585   ParsingOpenACCDirectiveRAII DirScope(*this);
586   ConsumeAnnotationToken();
587 
588   ParseOpenACCDirective();
589 
590   return nullptr;
591 }
592 
593 // Parse OpenACC Directive on a Statement.
594 StmtResult Parser::ParseOpenACCDirectiveStmt() {
595   assert(Tok.is(tok::annot_pragma_openacc) && "expected OpenACC Start Token");
596 
597   ParsingOpenACCDirectiveRAII DirScope(*this);
598   ConsumeAnnotationToken();
599 
600   ParseOpenACCDirective();
601 
602   return StmtEmpty();
603 }
604