[clang] [OpenACC] Implement Atomic construct variants (PR #73015)
Erich Keane via cfe-commits
cfe-commits at lists.llvm.org
Tue Nov 21 08:58:45 PST 2023
https://github.com/erichkeane created https://github.com/llvm/llvm-project/pull/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.
>From e656fb03feeafa1997a4b93126063ad77ecf5b2c Mon Sep 17 00:00:00 2001
From: erichkeane <ekeane at nvidia.com>
Date: Tue, 21 Nov 2023 08:56:00 -0800
Subject: [PATCH] [OpenACC] Implement Atomic construct variants
`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.
---
.../clang/Basic/DiagnosticParseKinds.td | 3 ++
clang/include/clang/Basic/OpenACCKinds.h | 8 +++-
clang/lib/Parse/ParseOpenACC.cpp | 43 ++++++++++++++++++-
clang/test/ParserOpenACC/parse-constructs.c | 31 +++++++++++++
4 files changed, 83 insertions(+), 2 deletions(-)
diff --git a/clang/include/clang/Basic/DiagnosticParseKinds.td b/clang/include/clang/Basic/DiagnosticParseKinds.td
index 54b5ba6e6414b2d..1e307f481d6c9ad 100644
--- a/clang/include/clang/Basic/DiagnosticParseKinds.td
+++ b/clang/include/clang/Basic/DiagnosticParseKinds.td
@@ -1364,6 +1364,9 @@ def warn_pragma_acc_unimplemented_clause_parsing
def err_acc_invalid_directive
: Error<"invalid OpenACC directive '%select{%1|%1 %2}0'">;
def err_acc_missing_directive : Error<"expected OpenACC directive">;
+def err_acc_invalid_atomic_clause
+ : Error<"%select{missing|invalid}0 OpenACC 'atomic-clause'%select{| "
+ "'%1'}0; expected 'read', 'write', 'update', or 'capture'">;
// OpenMP support.
def warn_pragma_omp_ignored : Warning<
diff --git a/clang/include/clang/Basic/OpenACCKinds.h b/clang/include/clang/Basic/OpenACCKinds.h
index 2a818638720abb0..1622f8f00274ad6 100644
--- a/clang/include/clang/Basic/OpenACCKinds.h
+++ b/clang/include/clang/Basic/OpenACCKinds.h
@@ -41,7 +41,13 @@ enum class OpenACCDirectiveKind {
SerialLoop,
KernelsLoop,
- // FIXME: atomic Construct variants.
+ // Atomic Construct. The OpenACC standard considers these as a single
+ // construct, however the atomic-clause (read, write, update, capture) are
+ // important for legalization of the application of this to statements/blocks.
+ AtomicRead,
+ AtomicWrite,
+ AtomicUpdate,
+ AtomicCapture,
// Declare Directive.
Declare,
diff --git a/clang/lib/Parse/ParseOpenACC.cpp b/clang/lib/Parse/ParseOpenACC.cpp
index a0f8fa97f6fa701..8a4f7a5da636e30 100644
--- a/clang/lib/Parse/ParseOpenACC.cpp
+++ b/clang/lib/Parse/ParseOpenACC.cpp
@@ -29,7 +29,8 @@ enum class OpenACCDirectiveKindEx {
// 'enter data' and 'exit data'
Enter,
Exit,
- // FIXME: Atomic Variants
+ // 'atomic read', 'atomic write', 'atomic update', and 'atomic capture'.
+ Atomic,
};
// Translate single-token string representations to the OpenACC Directive Kind.
@@ -59,9 +60,21 @@ OpenACCDirectiveKindEx getOpenACCDirectiveKind(StringRef Name) {
return llvm::StringSwitch<OpenACCDirectiveKindEx>(Name)
.Case("enter", OpenACCDirectiveKindEx::Enter)
.Case("exit", OpenACCDirectiveKindEx::Exit)
+ .Case("atomic", OpenACCDirectiveKindEx::Atomic)
.Default(OpenACCDirectiveKindEx::Invalid);
}
+// Since 'atomic' is effectively a compound directive, this will decode the
+// second part of the directive.
+OpenACCDirectiveKind getOpenACCAtomicDirectiveKind(StringRef Name) {
+ return llvm::StringSwitch<OpenACCDirectiveKind>(Name)
+ .Case("read", OpenACCDirectiveKind::AtomicRead)
+ .Case("write", OpenACCDirectiveKind::AtomicWrite)
+ .Case("update", OpenACCDirectiveKind::AtomicUpdate)
+ .Case("capture", OpenACCDirectiveKind::AtomicCapture)
+ .Default(OpenACCDirectiveKind::Invalid);
+}
+
bool isOpenACCDirectiveKind(OpenACCDirectiveKind Kind, StringRef Tok) {
switch (Kind) {
case OpenACCDirectiveKind::Parallel:
@@ -82,6 +95,10 @@ bool isOpenACCDirectiveKind(OpenACCDirectiveKind Kind, StringRef Tok) {
case OpenACCDirectiveKind::KernelsLoop:
case OpenACCDirectiveKind::EnterData:
case OpenACCDirectiveKind::ExitData:
+ case OpenACCDirectiveKind::AtomicRead:
+ case OpenACCDirectiveKind::AtomicWrite:
+ case OpenACCDirectiveKind::AtomicUpdate:
+ case OpenACCDirectiveKind::AtomicCapture:
return false;
case OpenACCDirectiveKind::Declare:
@@ -126,6 +143,28 @@ ParseOpenACCEnterExitDataDirective(Parser &P, Token FirstTok,
: OpenACCDirectiveKind::ExitData;
}
+OpenACCDirectiveKind ParseOpenACCAtomicDirective(Parser &P) {
+ Token AtomicClauseToken = P.getCurToken();
+
+ if (AtomicClauseToken.isAnnotation()) {
+ P.Diag(AtomicClauseToken, diag::err_acc_invalid_atomic_clause) << 0;
+ return OpenACCDirectiveKind::Invalid;
+ }
+
+ std::string AtomicClauseSpelling =
+ P.getPreprocessor().getSpelling(AtomicClauseToken);
+
+ OpenACCDirectiveKind DirKind =
+ getOpenACCAtomicDirectiveKind(AtomicClauseSpelling);
+
+ if (DirKind == OpenACCDirectiveKind::Invalid)
+ P.Diag(AtomicClauseToken, diag::err_acc_invalid_atomic_clause)
+ << 1 << AtomicClauseSpelling;
+
+ P.ConsumeToken();
+ return DirKind;
+}
+
// Parse and consume the tokens for OpenACC Directive/Construct kinds.
OpenACCDirectiveKind ParseOpenACCDirectiveKind(Parser &P) {
Token FirstTok = P.getCurToken();
@@ -158,6 +197,8 @@ OpenACCDirectiveKind ParseOpenACCDirectiveKind(Parser &P) {
case OpenACCDirectiveKindEx::Exit:
return ParseOpenACCEnterExitDataDirective(P, FirstTok, FirstTokSpelling,
ExDirKind);
+ case OpenACCDirectiveKindEx::Atomic:
+ return ParseOpenACCAtomicDirective(P);
}
}
diff --git a/clang/test/ParserOpenACC/parse-constructs.c b/clang/test/ParserOpenACC/parse-constructs.c
index a5270daf6034cf8..28410cff7f15d44 100644
--- a/clang/test/ParserOpenACC/parse-constructs.c
+++ b/clang/test/ParserOpenACC/parse-constructs.c
@@ -94,6 +94,37 @@ void func() {
#pragma acc kernels loop
for(;;){}
+ int i = 0, j = 0, k = 0;
+ // expected-error at +2{{missing OpenACC 'atomic-clause'; expected 'read', 'write', 'update', or 'capture'}}
+ // expected-warning at +1{{OpenACC directives not yet implemented, pragma ignored}}
+#pragma acc atomic
+ i = j;
+ // expected-error at +2{{invalid OpenACC 'atomic-clause' 'garbage'; expected 'read', 'write', 'update', or 'capture'}}
+ // expected-warning at +1{{OpenACC directives not yet implemented, pragma ignored}}
+#pragma acc atomic garbage
+ i = j;
+ // expected-warning at +3{{OpenACC clause parsing not yet implemented}}
+ // expected-error at +2{{invalid OpenACC 'atomic-clause' 'garbage'; expected 'read', 'write', 'update', or 'capture'}}
+ // expected-warning at +1{{OpenACC directives not yet implemented, pragma ignored}}
+#pragma acc atomic garbage clause list
+ i = j;
+ // expected-warning at +1{{OpenACC directives not yet implemented, pragma ignored}}
+#pragma acc atomic read
+ i = j;
+ // expected-warning at +2{{OpenACC clause parsing not yet implemented}}
+ // expected-warning at +1{{OpenACC directives not yet implemented, pragma ignored}}
+#pragma acc atomic write clause list
+ i = i + j;
+ // expected-warning at +2{{OpenACC clause parsing not yet implemented}}
+ // expected-warning at +1{{OpenACC directives not yet implemented, pragma ignored}}
+#pragma acc atomic update clause list
+ i++;
+ // expected-warning at +2{{OpenACC clause parsing not yet implemented}}
+ // expected-warning at +1{{OpenACC directives not yet implemented, pragma ignored}}
+#pragma acc atomic capture clause list
+ i = j++;
+
+
// expected-warning at +2{{OpenACC clause parsing not yet implemented}}
// expected-warning at +1{{OpenACC directives not yet implemented, pragma ignored}}
#pragma acc declare clause list
More information about the cfe-commits
mailing list