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