[OpenACC] Implement Atomic construct variants (#73015)

`atomic` is required to be followed by a special `atomic clause`, so
this patch manages the parsing of that. We are representing each of the
variants of the atomic construct as separate kinds, because they have
distinct rules/application/etc, and this should make it easier to check
rules in the future.

Author: erichkeane

clang/include/clang/Basic/OpenACCKinds.h

@@ -41,7 +41,8 @@ enum class OpenACCDirectiveKind {
   SerialLoop,
   KernelsLoop,
 
-  // FIXME: atomic Construct variants.
+  // Atomic Construct.
+  Atomic,
 
   // Declare Directive.
   Declare,
@@ -59,6 +60,14 @@ enum class OpenACCDirectiveKind {
   // Invalid.
   Invalid,
 };
+
+enum class OpenACCAtomicKind {
+  Read,
+  Write,
+  Update,
+  Capture,
+  Invalid,
+};
 } // namespace clang
 
 #endif // LLVM_CLANG_BASIC_OPENACCKINDS_H

clang/lib/Parse/ParseOpenACC.cpp

@@ -29,7 +29,6 @@ enum class OpenACCDirectiveKindEx {
   // 'enter data' and 'exit data'
   Enter,
   Exit,
-  // FIXME: Atomic Variants
 };
 
 // Translate single-token string representations to the OpenACC Directive Kind.
@@ -46,6 +45,7 @@ OpenACCDirectiveKindEx getOpenACCDirectiveKind(StringRef Name) {
           .Case("data", OpenACCDirectiveKind::Data)
           .Case("host_data", OpenACCDirectiveKind::HostData)
           .Case("loop", OpenACCDirectiveKind::Loop)
+          .Case("atomic", OpenACCDirectiveKind::Atomic)
           .Case("declare", OpenACCDirectiveKind::Declare)
           .Case("init", OpenACCDirectiveKind::Init)
           .Case("shutdown", OpenACCDirectiveKind::Shutdown)
@@ -62,6 +62,17 @@ OpenACCDirectiveKindEx getOpenACCDirectiveKind(StringRef Name) {
       .Default(OpenACCDirectiveKindEx::Invalid);
 }
 
+// Since 'atomic' is effectively a compound directive, this will decode the
+// second part of the directive.
+OpenACCAtomicKind getOpenACCAtomicKind(StringRef Name) {
+  return llvm::StringSwitch<OpenACCAtomicKind>(Name)
+      .Case("read", OpenACCAtomicKind::Read)
+      .Case("write", OpenACCAtomicKind::Write)
+      .Case("update", OpenACCAtomicKind::Update)
+      .Case("capture", OpenACCAtomicKind::Capture)
+      .Default(OpenACCAtomicKind::Invalid);
+}
+
 bool isOpenACCDirectiveKind(OpenACCDirectiveKind Kind, StringRef Tok) {
   switch (Kind) {
   case OpenACCDirectiveKind::Parallel:
@@ -84,6 +95,8 @@ bool isOpenACCDirectiveKind(OpenACCDirectiveKind Kind, StringRef Tok) {
   case OpenACCDirectiveKind::ExitData:
     return false;
 
+  case OpenACCDirectiveKind::Atomic:
+    return Tok == "atomic";
   case OpenACCDirectiveKind::Declare:
     return Tok == "declare";
   case OpenACCDirectiveKind::Init:
@@ -126,6 +139,27 @@ ParseOpenACCEnterExitDataDirective(Parser &P, Token FirstTok,
              : OpenACCDirectiveKind::ExitData;
 }
 
+OpenACCAtomicKind ParseOpenACCAtomicKind(Parser &P) {
+  Token AtomicClauseToken = P.getCurToken();
+
+  // #pragma acc atomic is equivilent to update:
+  if (AtomicClauseToken.isAnnotation())
+    return OpenACCAtomicKind::Update;
+
+  std::string AtomicClauseSpelling =
+      P.getPreprocessor().getSpelling(AtomicClauseToken);
+  OpenACCAtomicKind AtomicKind = getOpenACCAtomicKind(AtomicClauseSpelling);
+
+  // If we don't know what this is, treat it as 'nothing', and treat the rest of
+  // this as a clause list, which, despite being invalid, is likely what the
+  // user was trying to do.
+  if (AtomicKind == OpenACCAtomicKind::Invalid)
+    return OpenACCAtomicKind::Update;
+
+  P.ConsumeToken();
+  return AtomicKind;
+}
+
 // Parse and consume the tokens for OpenACC Directive/Construct kinds.
 OpenACCDirectiveKind ParseOpenACCDirectiveKind(Parser &P) {
   Token FirstTok = P.getCurToken();
@@ -199,7 +233,13 @@ void ParseOpenACCClauseList(Parser &P) {
 }
 
 void ParseOpenACCDirective(Parser &P) {
-  ParseOpenACCDirectiveKind(P);
+  OpenACCDirectiveKind DirKind = ParseOpenACCDirectiveKind(P);
+
+  // Once we've parsed the construct/directive name, some have additional
+  // specifiers that need to be taken care of. Atomic has an 'atomic-clause'
+  // that needs to be parsed.
+  if (DirKind == OpenACCDirectiveKind::Atomic)
+    ParseOpenACCAtomicKind(P);
 
   // Parses the list of clauses, if present.
   ParseOpenACCClauseList(P);

clang/test/ParserOpenACC/parse-constructs.c

@@ -94,6 +94,35 @@ void func() {
 #pragma acc kernels loop
   for(;;){}
 
+  int i = 0, j = 0, k = 0;
+  // expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
+#pragma acc atomic
+  i = j;
+  // expected-warning@+2{{OpenACC clause parsing not yet implemented}}
+  // expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
+#pragma acc atomic garbage
+  i = j;
+  // expected-warning@+2{{OpenACC clause parsing not yet implemented}}
+  // expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
+#pragma acc atomic garbage clause list
+  i = j;
+  // expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
+#pragma acc atomic read
+  i = j;
+  // expected-warning@+2{{OpenACC clause parsing not yet implemented}}
+  // expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
+#pragma acc atomic write clause list
+  i = i + j;
+  // expected-warning@+2{{OpenACC clause parsing not yet implemented}}
+  // expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
+#pragma acc atomic update clause list
+  i++;
+  // expected-warning@+2{{OpenACC clause parsing not yet implemented}}
+  // expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
+#pragma acc atomic capture clause list
+  i = j++;
+
+
   // expected-warning@+2{{OpenACC clause parsing not yet implemented}}
   // expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
 #pragma acc declare clause list