[clang] [OpenACC] Implement Atomic construct variants (PR #73015)
Erich Keane via cfe-commits
cfe-commits at lists.llvm.org
Tue Nov 21 09:59:50 PST 2023
https://github.com/erichkeane updated https://github.com/llvm/llvm-project/pull/73015
>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 1/2] [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
>From 3a3450a176e42727336a57daf4d76cb92b94d07d Mon Sep 17 00:00:00 2001
From: erichkeane <ekeane at nvidia.com>
Date: Tue, 21 Nov 2023 09:57:45 -0800
Subject: [PATCH 2/2] Refactor to no longer consider atomic variants as
constructs
After review, it was pointed out that this is a list that has changed
recently in OpenMP and thus might change here.
This patch extracts the parsing of atomic-clause to happen after
directive determination.
---
.../clang/Basic/DiagnosticParseKinds.td | 3 -
clang/include/clang/Basic/OpenACCKinds.h | 17 +++---
clang/lib/Parse/ParseOpenACC.cpp | 57 +++++++++----------
clang/test/ParserOpenACC/parse-constructs.c | 6 +-
4 files changed, 40 insertions(+), 43 deletions(-)
diff --git a/clang/include/clang/Basic/DiagnosticParseKinds.td b/clang/include/clang/Basic/DiagnosticParseKinds.td
index 1e307f481d6c9ad..54b5ba6e6414b2d 100644
--- a/clang/include/clang/Basic/DiagnosticParseKinds.td
+++ b/clang/include/clang/Basic/DiagnosticParseKinds.td
@@ -1364,9 +1364,6 @@ 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 1622f8f00274ad6..cf4bad9ce0cb9ff 100644
--- a/clang/include/clang/Basic/OpenACCKinds.h
+++ b/clang/include/clang/Basic/OpenACCKinds.h
@@ -41,13 +41,8 @@ enum class OpenACCDirectiveKind {
SerialLoop,
KernelsLoop,
- // 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,
+ // Atomic Construct.
+ Atomic,
// Declare Directive.
Declare,
@@ -65,6 +60,14 @@ enum class OpenACCDirectiveKind {
// Invalid.
Invalid,
};
+
+enum class OpenACCAtomicKind {
+ Read,
+ Write,
+ Update,
+ Capture,
+ Invalid,
+};
} // namespace clang
#endif // LLVM_CLANG_BASIC_OPENACCKINDS_H
diff --git a/clang/lib/Parse/ParseOpenACC.cpp b/clang/lib/Parse/ParseOpenACC.cpp
index 8a4f7a5da636e30..978a07ec82e4288 100644
--- a/clang/lib/Parse/ParseOpenACC.cpp
+++ b/clang/lib/Parse/ParseOpenACC.cpp
@@ -29,8 +29,6 @@ enum class OpenACCDirectiveKindEx {
// 'enter data' and 'exit data'
Enter,
Exit,
- // 'atomic read', 'atomic write', 'atomic update', and 'atomic capture'.
- Atomic,
};
// Translate single-token string representations to the OpenACC Directive Kind.
@@ -47,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)
@@ -60,19 +59,18 @@ 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);
+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) {
@@ -95,12 +93,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::Atomic:
+ return Tok == "atomic";
case OpenACCDirectiveKind::Declare:
return Tok == "declare";
case OpenACCDirectiveKind::Init:
@@ -143,26 +139,25 @@ ParseOpenACCEnterExitDataDirective(Parser &P, Token FirstTok,
: OpenACCDirectiveKind::ExitData;
}
-OpenACCDirectiveKind ParseOpenACCAtomicDirective(Parser &P) {
+OpenACCAtomicKind ParseOpenACCAtomicKind(Parser &P) {
Token AtomicClauseToken = P.getCurToken();
- if (AtomicClauseToken.isAnnotation()) {
- P.Diag(AtomicClauseToken, diag::err_acc_invalid_atomic_clause) << 0;
- return OpenACCDirectiveKind::Invalid;
- }
+ // #pragma acc atomic is equivilent to update:
+ if (AtomicClauseToken.isAnnotation())
+ return OpenACCAtomicKind::Update;
std::string AtomicClauseSpelling =
P.getPreprocessor().getSpelling(AtomicClauseToken);
+ OpenACCAtomicKind AtomicKind = getOpenACCAtomicKind(AtomicClauseSpelling);
- OpenACCDirectiveKind DirKind =
- getOpenACCAtomicDirectiveKind(AtomicClauseSpelling);
-
- if (DirKind == OpenACCDirectiveKind::Invalid)
- P.Diag(AtomicClauseToken, diag::err_acc_invalid_atomic_clause)
- << 1 << 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 DirKind;
+ return AtomicKind;
}
// Parse and consume the tokens for OpenACC Directive/Construct kinds.
@@ -197,8 +192,6 @@ OpenACCDirectiveKind ParseOpenACCDirectiveKind(Parser &P) {
case OpenACCDirectiveKindEx::Exit:
return ParseOpenACCEnterExitDataDirective(P, FirstTok, FirstTokSpelling,
ExDirKind);
- case OpenACCDirectiveKindEx::Atomic:
- return ParseOpenACCAtomicDirective(P);
}
}
@@ -240,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);
diff --git a/clang/test/ParserOpenACC/parse-constructs.c b/clang/test/ParserOpenACC/parse-constructs.c
index 28410cff7f15d44..59d14cff9d416e9 100644
--- a/clang/test/ParserOpenACC/parse-constructs.c
+++ b/clang/test/ParserOpenACC/parse-constructs.c
@@ -95,16 +95,14 @@ void func() {
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 +2{{OpenACC clause parsing not yet implemented}}
// 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 +2{{OpenACC clause parsing not yet implemented}}
// expected-warning at +1{{OpenACC directives not yet implemented, pragma ignored}}
#pragma acc atomic garbage clause list
i = j;
More information about the cfe-commits
mailing list