[clang] 99a9133 - [OpenACC] Implement Sema/AST for 'atomic' construct
via cfe-commits
cfe-commits at lists.llvm.org
Mon Feb 3 07:24:03 PST 2025
Author: erichkeane
Date: 2025-02-03T07:22:22-08:00
New Revision: 99a9133a68b77cb978dd4b0cdbcd67e4edf7bd92
URL: https://github.com/llvm/llvm-project/commit/99a9133a68b77cb978dd4b0cdbcd67e4edf7bd92
DIFF: https://github.com/llvm/llvm-project/commit/99a9133a68b77cb978dd4b0cdbcd67e4edf7bd92.diff
LOG: [OpenACC] Implement Sema/AST for 'atomic' construct
The atomic construct is a particularly complicated one. The directive
itself is pretty simple, it has 5 options for the 'atomic-clause'.
However, the associated statement is fairly complicated.
'read' accepts:
v = x;
'write' accepts:
x = expr;
'update' (or no clause) accepts:
x++;
x--;
++x;
--x;
x binop= expr;
x = x binop expr;
x = expr binop x;
'capture' accepts either a compound statement, or:
v = x++;
v = x--;
v = ++x;
v = --x;
v = x binop= expr;
v = x = x binop expr;
v = x = expr binop x;
IF 'capture' has a compound statement, it accepts:
{v = x; x binop= expr; }
{x binop= expr; v = x; }
{v = x; x = x binop expr; }
{v = x; x = expr binop x; }
{x = x binop expr ;v = x; }
{x = expr binop x; v = x; }
{v = x; x = expr; }
{v = x; x++; }
{v = x; ++x; }
{x++; v = x; }
{++x; v = x; }
{v = x; x--; }
{v = x; --x; }
{x--; v = x; }
{--x; v = x; }
While these are all quite complicated, there is a significant amount
of similarity between the 'capture' and 'update' lists, so this patch
reuses a lot of the same functions.
This patch implements the entirety of 'atomic', creating a new Sema file
for the sema for it, as it is fairly sizable.
Added:
clang/lib/Sema/SemaOpenACCAtomic.cpp
clang/test/AST/ast-print-openacc-atomic-construct.cpp
clang/test/SemaOpenACC/atomic-construct-ast.cpp
clang/test/SemaOpenACC/atomic-construct.cpp
Modified:
clang/include/clang-c/Index.h
clang/include/clang/AST/RecursiveASTVisitor.h
clang/include/clang/AST/StmtOpenACC.h
clang/include/clang/AST/TextNodeDumper.h
clang/include/clang/Basic/DiagnosticSemaKinds.td
clang/include/clang/Basic/OpenACCKinds.h
clang/include/clang/Basic/StmtNodes.td
clang/include/clang/Parse/Parser.h
clang/include/clang/Sema/SemaOpenACC.h
clang/include/clang/Serialization/ASTBitCodes.h
clang/lib/AST/StmtOpenACC.cpp
clang/lib/AST/StmtPrinter.cpp
clang/lib/AST/StmtProfile.cpp
clang/lib/AST/TextNodeDumper.cpp
clang/lib/CodeGen/CGStmt.cpp
clang/lib/CodeGen/CodeGenFunction.h
clang/lib/Parse/ParseOpenACC.cpp
clang/lib/Sema/CMakeLists.txt
clang/lib/Sema/SemaExceptionSpec.cpp
clang/lib/Sema/SemaOpenACC.cpp
clang/lib/Sema/SemaOpenACCClause.cpp
clang/lib/Sema/TreeTransform.h
clang/lib/Serialization/ASTReaderStmt.cpp
clang/lib/Serialization/ASTWriterStmt.cpp
clang/lib/StaticAnalyzer/Core/ExprEngine.cpp
clang/test/ParserOpenACC/parse-constructs.c
clang/tools/libclang/CIndex.cpp
clang/tools/libclang/CXCursor.cpp
Removed:
################################################################################
diff --git a/clang/include/clang-c/Index.h b/clang/include/clang-c/Index.h
index cc7c65b15088476..61e361faabdaf8b 100644
--- a/clang/include/clang-c/Index.h
+++ b/clang/include/clang-c/Index.h
@@ -2206,7 +2206,11 @@ enum CXCursorKind {
*/
CXCursor_OpenACCUpdateConstruct = 331,
- CXCursor_LastStmt = CXCursor_OpenACCUpdateConstruct,
+ /** OpenACC atomic Construct.
+ */
+ CXCursor_OpenACCAtomicConstruct = 332,
+
+ CXCursor_LastStmt = CXCursor_OpenACCAtomicConstruct,
/**
* Cursor that represents the translation unit itself.
diff --git a/clang/include/clang/AST/RecursiveASTVisitor.h b/clang/include/clang/AST/RecursiveASTVisitor.h
index 5f4c39b9cbdb75d..55505794e70542c 100644
--- a/clang/include/clang/AST/RecursiveASTVisitor.h
+++ b/clang/include/clang/AST/RecursiveASTVisitor.h
@@ -4099,6 +4099,8 @@ DEF_TRAVERSE_STMT(OpenACCSetConstruct,
{ TRY_TO(VisitOpenACCClauseList(S->clauses())); })
DEF_TRAVERSE_STMT(OpenACCUpdateConstruct,
{ TRY_TO(VisitOpenACCClauseList(S->clauses())); })
+DEF_TRAVERSE_STMT(OpenACCAtomicConstruct,
+ { TRY_TO(TraverseOpenACCAssociatedStmtConstruct(S)); })
// Traverse HLSL: Out argument expression
DEF_TRAVERSE_STMT(HLSLOutArgExpr, {})
diff --git a/clang/include/clang/AST/StmtOpenACC.h b/clang/include/clang/AST/StmtOpenACC.h
index ebbee152f918f82..bd6c95d342ce2ec 100644
--- a/clang/include/clang/AST/StmtOpenACC.h
+++ b/clang/include/clang/AST/StmtOpenACC.h
@@ -751,5 +751,50 @@ class OpenACCUpdateConstruct final
Create(const ASTContext &C, SourceLocation Start, SourceLocation DirectiveLoc,
SourceLocation End, ArrayRef<const OpenACCClause *> Clauses);
};
+
+// This class represents the 'atomic' construct, which has an associated
+// statement, but no clauses.
+class OpenACCAtomicConstruct final : public OpenACCAssociatedStmtConstruct {
+
+ friend class ASTStmtReader;
+ OpenACCAtomicKind AtomicKind = OpenACCAtomicKind::None;
+
+ OpenACCAtomicConstruct(EmptyShell)
+ : OpenACCAssociatedStmtConstruct(
+ OpenACCAtomicConstructClass, OpenACCDirectiveKind::Atomic,
+ SourceLocation{}, SourceLocation{}, SourceLocation{},
+ /*AssociatedStmt=*/nullptr) {}
+
+ OpenACCAtomicConstruct(SourceLocation Start, SourceLocation DirectiveLoc,
+ OpenACCAtomicKind AtKind, SourceLocation End,
+ Stmt *AssociatedStmt)
+ : OpenACCAssociatedStmtConstruct(OpenACCAtomicConstructClass,
+ OpenACCDirectiveKind::Atomic, Start,
+ DirectiveLoc, End, AssociatedStmt),
+ AtomicKind(AtKind) {}
+
+ void setAssociatedStmt(Stmt *S) {
+ OpenACCAssociatedStmtConstruct::setAssociatedStmt(S);
+ }
+
+public:
+ static bool classof(const Stmt *T) {
+ return T->getStmtClass() == OpenACCAtomicConstructClass;
+ }
+
+ static OpenACCAtomicConstruct *CreateEmpty(const ASTContext &C);
+ static OpenACCAtomicConstruct *
+ Create(const ASTContext &C, SourceLocation Start, SourceLocation DirectiveLoc,
+ OpenACCAtomicKind AtKind, SourceLocation End, Stmt *AssociatedStmt);
+
+ OpenACCAtomicKind getAtomicKind() const { return AtomicKind; }
+ const Stmt *getAssociatedStmt() const {
+ return OpenACCAssociatedStmtConstruct::getAssociatedStmt();
+ }
+ Stmt *getAssociatedStmt() {
+ return OpenACCAssociatedStmtConstruct::getAssociatedStmt();
+ }
+};
+
} // namespace clang
#endif // LLVM_CLANG_AST_STMTOPENACC_H
diff --git a/clang/include/clang/AST/TextNodeDumper.h b/clang/include/clang/AST/TextNodeDumper.h
index 4aaae48ba8b42fa..bfd205ffb0d99ad 100644
--- a/clang/include/clang/AST/TextNodeDumper.h
+++ b/clang/include/clang/AST/TextNodeDumper.h
@@ -420,6 +420,7 @@ class TextNodeDumper
void VisitOpenACCSetConstruct(const OpenACCSetConstruct *S);
void VisitOpenACCShutdownConstruct(const OpenACCShutdownConstruct *S);
void VisitOpenACCUpdateConstruct(const OpenACCUpdateConstruct *S);
+ void VisitOpenACCAtomicConstruct(const OpenACCAtomicConstruct *S);
void VisitOpenACCAsteriskSizeExpr(const OpenACCAsteriskSizeExpr *S);
void VisitEmbedExpr(const EmbedExpr *S);
void VisitAtomicExpr(const AtomicExpr *AE);
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index 00a94eb7a303671..8bacb1b73459f14 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -12903,6 +12903,48 @@ def err_acc_update_as_body
: Error<"OpenACC 'update' construct may not appear in place of the "
"statement following a%select{n if statement| while statement| do "
"statement| switch statement| label statement}0">;
+def err_acc_invalid_atomic
+ : Error<"statement associated with OpenACC 'atomic%select{| "
+ "%1}0' directive is invalid">;
+def note_acc_atomic_expr_must_be
+ : Note<"expected "
+ "%enum_select<OACCAtomicExpr>{%Assign{assignment}|%UnaryCompAssign{"
+ "assignment, compound assignment, increment, or decrement}}0 "
+ "expression">;
+def note_acc_atomic_unsupported_unary_operator
+ : Note<"unary operator not supported, only increment and decrement "
+ "operations permitted">;
+def note_acc_atomic_unsupported_binary_operator
+ : Note<"binary operator not supported, only +, *, -, /, &, ^, |, <<, or >> "
+ "are permitted">;
+def note_acc_atomic_unsupported_compound_binary_operator
+ : Note<"compound binary operator not supported, only +=, *=, -=, /=, &=, "
+ "^=, |=, <<=, or >>= are permitted">;
+
+def note_acc_atomic_operand_lvalue_scalar
+ : Note<"%select{left |right |}0operand to "
+ "%enum_select<OACCAtomicOpKind>{%Assign{assignment}|%CompoundAssign{"
+ "compound assignment}|%Inc{increment}|"
+ "%Dec{decrement}}1 "
+ "expression must be "
+ "%enum_select<OACCLValScalar>{%LVal{an l-value}|%Scalar{of scalar "
+ "type (was %3)}}2">;
+def note_acc_atomic_too_many_stmts
+ : Note<"'atomic capture' with a compound statement only supports two "
+ "statements">;
+def note_acc_atomic_expected_binop : Note<"expected binary operation on right "
+ "hand side of assignment operator">;
+def note_acc_atomic_mismatch_operand
+ : Note<"left hand side of assignment operation('%0') must match one side "
+ "of the sub-operation on the right hand side('%1' and '%2')">;
+def note_acc_atomic_mismatch_compound_operand
+ : Note<"variable %select{|in unary expression|on right hand side of "
+ "assignment|on left hand side of assignment|on left hand side of "
+ "compound assignment|on left hand side of assignment}2('%3') must "
+ "match variable used %select{|in unary expression|on right hand "
+ "side of assignment|<not possible>|on left hand side of compound "
+ "assignment|on left hand side of assignment}0('%1') from the first "
+ "statement">;
// AMDGCN builtins diagnostics
def err_amdgcn_global_load_lds_size_invalid_value : Error<"invalid size value">;
diff --git a/clang/include/clang/Basic/OpenACCKinds.h b/clang/include/clang/Basic/OpenACCKinds.h
index 7fb76271826a6b0..739422caad64592 100644
--- a/clang/include/clang/Basic/OpenACCKinds.h
+++ b/clang/include/clang/Basic/OpenACCKinds.h
@@ -171,9 +171,33 @@ enum class OpenACCAtomicKind : uint8_t {
Write,
Update,
Capture,
- Invalid,
+ None,
};
+template <typename StreamTy>
+inline StreamTy &printOpenACCAtomicKind(StreamTy &Out, OpenACCAtomicKind AK) {
+ switch (AK) {
+ case OpenACCAtomicKind::Read:
+ return Out << "read";
+ case OpenACCAtomicKind::Write:
+ return Out << "write";
+ case OpenACCAtomicKind::Update:
+ return Out << "update";
+ case OpenACCAtomicKind::Capture:
+ return Out << "capture";
+ case OpenACCAtomicKind::None:
+ return Out << "<none>";
+ }
+}
+inline const StreamingDiagnostic &operator<<(const StreamingDiagnostic &Out,
+ OpenACCAtomicKind AK) {
+ return printOpenACCAtomicKind(Out, AK);
+}
+inline llvm::raw_ostream &operator<<(llvm::raw_ostream &Out,
+ OpenACCAtomicKind AK) {
+ return printOpenACCAtomicKind(Out, AK);
+}
+
/// Represents the kind of an OpenACC clause.
enum class OpenACCClauseKind : uint8_t {
/// 'finalize' clause, allowed on 'exit data' directive.
diff --git a/clang/include/clang/Basic/StmtNodes.td b/clang/include/clang/Basic/StmtNodes.td
index 2fea05e322c7547..d47e0a8157fc683 100644
--- a/clang/include/clang/Basic/StmtNodes.td
+++ b/clang/include/clang/Basic/StmtNodes.td
@@ -319,6 +319,7 @@ def OpenACCInitConstruct : StmtNode<OpenACCConstructStmt>;
def OpenACCShutdownConstruct : StmtNode<OpenACCConstructStmt>;
def OpenACCSetConstruct : StmtNode<OpenACCConstructStmt>;
def OpenACCUpdateConstruct : StmtNode<OpenACCConstructStmt>;
+def OpenACCAtomicConstruct : StmtNode<OpenACCAssociatedStmtConstruct>;
// OpenACC Additional Expressions.
def OpenACCAsteriskSizeExpr : StmtNode<Expr>;
diff --git a/clang/include/clang/Parse/Parser.h b/clang/include/clang/Parse/Parser.h
index e99d2cf2eaa4096..335258d597028f4 100644
--- a/clang/include/clang/Parse/Parser.h
+++ b/clang/include/clang/Parse/Parser.h
@@ -3710,6 +3710,7 @@ class Parser : public CodeCompletionHandler {
SourceLocation RParenLoc;
SourceLocation EndLoc;
SourceLocation MiscLoc;
+ OpenACCAtomicKind AtomicKind;
SmallVector<Expr *> Exprs;
SmallVector<OpenACCClause *> Clauses;
// TODO OpenACC: As we implement support for the Atomic, Routine, and Cache
diff --git a/clang/include/clang/Sema/SemaOpenACC.h b/clang/include/clang/Sema/SemaOpenACC.h
index 2e5a0ea0aaac64a..3004b98760a98eb 100644
--- a/clang/include/clang/Sema/SemaOpenACC.h
+++ b/clang/include/clang/Sema/SemaOpenACC.h
@@ -695,24 +695,53 @@ class SemaOpenACC : public SemaBase {
/// should check legality of the statement as it appertains to this Construct.
StmtResult ActOnAssociatedStmt(SourceLocation DirectiveLoc,
OpenACCDirectiveKind K,
+ OpenACCAtomicKind AtKind,
ArrayRef<const OpenACCClause *> Clauses,
StmtResult AssocStmt);
+ StmtResult ActOnAssociatedStmt(SourceLocation DirectiveLoc,
+ OpenACCDirectiveKind K,
+ ArrayRef<const OpenACCClause *> Clauses,
+ StmtResult AssocStmt) {
+ return ActOnAssociatedStmt(DirectiveLoc, K, OpenACCAtomicKind::None,
+ Clauses, AssocStmt);
+ }
+ /// Called to check the form of the `atomic` construct which has some fairly
+ /// sizable restrictions.
+ StmtResult CheckAtomicAssociatedStmt(SourceLocation AtomicDirLoc,
+ OpenACCAtomicKind AtKind,
+ StmtResult AssocStmt);
+
/// Called after the directive has been completely parsed, including the
/// declaration group or associated statement.
+ /// DirLoc: Location of the actual directive keyword.
/// LParenLoc: Location of the left paren, if it exists (not on all
/// constructs).
/// MiscLoc: First misc location, if necessary (not all constructs).
/// Exprs: List of expressions on the construct itself, if necessary (not all
/// constructs).
+ /// AK: The atomic kind of the directive, if necessary (atomic only)
/// RParenLoc: Location of the right paren, if it exists (not on all
/// constructs).
+ /// EndLoc: The last source location of the driective.
+ /// Clauses: The list of clauses for the directive, if present.
+ /// AssocStmt: The associated statement for this construct, if necessary.
StmtResult ActOnEndStmtDirective(
OpenACCDirectiveKind K, SourceLocation StartLoc, SourceLocation DirLoc,
SourceLocation LParenLoc, SourceLocation MiscLoc, ArrayRef<Expr *> Exprs,
- SourceLocation RParenLoc, SourceLocation EndLoc,
+ OpenACCAtomicKind AK, SourceLocation RParenLoc, SourceLocation EndLoc,
ArrayRef<OpenACCClause *> Clauses, StmtResult AssocStmt);
+ StmtResult ActOnEndStmtDirective(
+ OpenACCDirectiveKind K, SourceLocation StartLoc, SourceLocation DirLoc,
+ SourceLocation LParenLoc, SourceLocation MiscLoc, ArrayRef<Expr *> Exprs,
+ SourceLocation RParenLoc, SourceLocation EndLoc,
+ ArrayRef<OpenACCClause *> Clauses, StmtResult AssocStmt) {
+ return ActOnEndStmtDirective(K, StartLoc, DirLoc, LParenLoc, MiscLoc, Exprs,
+ OpenACCAtomicKind::None, RParenLoc, EndLoc,
+ Clauses, AssocStmt);
+ }
+
/// Called after the directive has been completely parsed, including the
/// declaration group or associated statement.
DeclGroupRef ActOnEndDeclDirective();
diff --git a/clang/include/clang/Serialization/ASTBitCodes.h b/clang/include/clang/Serialization/ASTBitCodes.h
index 7656add0b6884cc..3c184db5b2adfb5 100644
--- a/clang/include/clang/Serialization/ASTBitCodes.h
+++ b/clang/include/clang/Serialization/ASTBitCodes.h
@@ -2045,6 +2045,7 @@ enum StmtCode {
STMT_OPENACC_SHUTDOWN_CONSTRUCT,
STMT_OPENACC_SET_CONSTRUCT,
STMT_OPENACC_UPDATE_CONSTRUCT,
+ STMT_OPENACC_ATOMIC_CONSTRUCT,
// HLSL Constructs
EXPR_HLSL_OUT_ARG,
diff --git a/clang/lib/AST/StmtOpenACC.cpp b/clang/lib/AST/StmtOpenACC.cpp
index 2b0ac716bab56f0..11eab0c27579d85 100644
--- a/clang/lib/AST/StmtOpenACC.cpp
+++ b/clang/lib/AST/StmtOpenACC.cpp
@@ -305,3 +305,19 @@ OpenACCUpdateConstruct::Create(const ASTContext &C, SourceLocation Start,
new (Mem) OpenACCUpdateConstruct(Start, DirectiveLoc, End, Clauses);
return Inst;
}
+
+OpenACCAtomicConstruct *
+OpenACCAtomicConstruct::CreateEmpty(const ASTContext &C) {
+ void *Mem = C.Allocate(sizeof(OpenACCAtomicConstruct));
+ auto *Inst = new (Mem) OpenACCAtomicConstruct(EmptyShell{});
+ return Inst;
+}
+
+OpenACCAtomicConstruct *OpenACCAtomicConstruct::Create(
+ const ASTContext &C, SourceLocation Start, SourceLocation DirectiveLoc,
+ OpenACCAtomicKind AtKind, SourceLocation End, Stmt *AssociatedStmt) {
+ void *Mem = C.Allocate(sizeof(OpenACCAtomicConstruct));
+ auto *Inst = new (Mem)
+ OpenACCAtomicConstruct(Start, DirectiveLoc, AtKind, End, AssociatedStmt);
+ return Inst;
+}
diff --git a/clang/lib/AST/StmtPrinter.cpp b/clang/lib/AST/StmtPrinter.cpp
index d523abfe312848f..bae4134e2835941 100644
--- a/clang/lib/AST/StmtPrinter.cpp
+++ b/clang/lib/AST/StmtPrinter.cpp
@@ -1242,6 +1242,16 @@ void StmtPrinter::VisitOpenACCWaitConstruct(OpenACCWaitConstruct *S) {
OS << '\n';
}
+void StmtPrinter::VisitOpenACCAtomicConstruct(OpenACCAtomicConstruct *S) {
+ Indent() << "#pragma acc atomic";
+
+ if (S->getAtomicKind() != OpenACCAtomicKind::None)
+ OS << " " << S->getAtomicKind();
+
+ OS << '\n';
+ PrintStmt(S->getAssociatedStmt());
+}
+
//===----------------------------------------------------------------------===//
// Expr printing methods.
//===----------------------------------------------------------------------===//
diff --git a/clang/lib/AST/StmtProfile.cpp b/clang/lib/AST/StmtProfile.cpp
index 84985fcb20ff9ee..36d231e21fa9509 100644
--- a/clang/lib/AST/StmtProfile.cpp
+++ b/clang/lib/AST/StmtProfile.cpp
@@ -2809,6 +2809,11 @@ void StmtProfiler::VisitOpenACCUpdateConstruct(
P.VisitOpenACCClauseList(S->clauses());
}
+void StmtProfiler::VisitOpenACCAtomicConstruct(
+ const OpenACCAtomicConstruct *S) {
+ VisitStmt(S);
+}
+
void StmtProfiler::VisitHLSLOutArgExpr(const HLSLOutArgExpr *S) {
VisitStmt(S);
}
diff --git a/clang/lib/AST/TextNodeDumper.cpp b/clang/lib/AST/TextNodeDumper.cpp
index bbb31a1e8d3f5c9..10d7e4c0c738721 100644
--- a/clang/lib/AST/TextNodeDumper.cpp
+++ b/clang/lib/AST/TextNodeDumper.cpp
@@ -3041,6 +3041,12 @@ void TextNodeDumper::VisitOpenACCUpdateConstruct(
VisitOpenACCConstructStmt(S);
}
+void TextNodeDumper::VisitOpenACCAtomicConstruct(
+ const OpenACCAtomicConstruct *S) {
+ VisitOpenACCConstructStmt(S);
+ OS << ' ' << S->getAtomicKind();
+}
+
void TextNodeDumper::VisitEmbedExpr(const EmbedExpr *S) {
AddChild("begin", [=] { OS << S->getStartingElementPos(); });
AddChild("number of elements", [=] { OS << S->getDataElementCount(); });
diff --git a/clang/lib/CodeGen/CGStmt.cpp b/clang/lib/CodeGen/CGStmt.cpp
index 7c944fe85a352d5..e2ae1046c084a8e 100644
--- a/clang/lib/CodeGen/CGStmt.cpp
+++ b/clang/lib/CodeGen/CGStmt.cpp
@@ -489,6 +489,8 @@ void CodeGenFunction::EmitStmt(const Stmt *S, ArrayRef<const Attr *> Attrs) {
case Stmt::OpenACCUpdateConstructClass:
EmitOpenACCUpdateConstruct(cast<OpenACCUpdateConstruct>(*S));
break;
+ case Stmt::OpenACCAtomicConstructClass:
+ EmitOpenACCAtomicConstruct(cast<OpenACCAtomicConstruct>(*S));
}
}
diff --git a/clang/lib/CodeGen/CodeGenFunction.h b/clang/lib/CodeGen/CodeGenFunction.h
index 670dfce27f6a26c..ced3484fbd2b6ce 100644
--- a/clang/lib/CodeGen/CodeGenFunction.h
+++ b/clang/lib/CodeGen/CodeGenFunction.h
@@ -4167,6 +4167,13 @@ class CodeGenFunction : public CodeGenTypeCache {
// but in the future we will implement some sort of IR.
}
+ void EmitOpenACCAtomicConstruct(const OpenACCAtomicConstruct &S) {
+ // TODO OpenACC: Implement this. It is currently implemented as a 'no-op',
+ // simply emitting its associated stmt, but in the future we will implement
+ // some sort of IR.
+ EmitStmt(S.getAssociatedStmt());
+ }
+
//===--------------------------------------------------------------------===//
// LValue Expression Emission
//===--------------------------------------------------------------------===//
diff --git a/clang/lib/Parse/ParseOpenACC.cpp b/clang/lib/Parse/ParseOpenACC.cpp
index 98fd61913e5a46e..d036971d2fc31c1 100644
--- a/clang/lib/Parse/ParseOpenACC.cpp
+++ b/clang/lib/Parse/ParseOpenACC.cpp
@@ -156,14 +156,14 @@ OpenACCClauseKind getOpenACCClauseKind(Token Tok) {
// second part of the directive.
OpenACCAtomicKind getOpenACCAtomicKind(Token Tok) {
if (!Tok.is(tok::identifier))
- return OpenACCAtomicKind::Invalid;
+ return OpenACCAtomicKind::None;
return llvm::StringSwitch<OpenACCAtomicKind>(
Tok.getIdentifierInfo()->getName())
.Case("read", OpenACCAtomicKind::Read)
.Case("write", OpenACCAtomicKind::Write)
.Case("update", OpenACCAtomicKind::Update)
.Case("capture", OpenACCAtomicKind::Capture)
- .Default(OpenACCAtomicKind::Invalid);
+ .Default(OpenACCAtomicKind::None);
}
OpenACCDefaultClauseKind getOpenACCDefaultClauseKind(Token Tok) {
@@ -398,17 +398,16 @@ OpenACCAtomicKind ParseOpenACCAtomicKind(Parser &P) {
// #pragma acc atomic is equivilent to update:
if (AtomicClauseToken.isAnnotation())
- return OpenACCAtomicKind::Update;
+ return OpenACCAtomicKind::None;
OpenACCAtomicKind AtomicKind = getOpenACCAtomicKind(AtomicClauseToken);
- // 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;
+ // If this isn't a valid atomic-kind, don't consume the token, and treat the
+ // rest as a clause list, which despite there being no permissible clauses,
+ // will diagnose as a clause.
+ if (AtomicKind != OpenACCAtomicKind::None)
+ P.ConsumeToken();
- P.ConsumeToken();
return AtomicKind;
}
@@ -570,12 +569,19 @@ void SkipUntilEndOfDirective(Parser &P) {
bool doesDirectiveHaveAssociatedStmt(OpenACCDirectiveKind DirKind) {
switch (DirKind) {
- default:
+ case OpenACCDirectiveKind::Routine:
+ // FIXME: Routine MIGHT end up needing to be 'true' here, as it needs a way
+ // to capture a lambda-expression on the next line.
+ case OpenACCDirectiveKind::Cache:
+ case OpenACCDirectiveKind::Declare:
+ case OpenACCDirectiveKind::Set:
case OpenACCDirectiveKind::EnterData:
case OpenACCDirectiveKind::ExitData:
case OpenACCDirectiveKind::Wait:
case OpenACCDirectiveKind::Init:
case OpenACCDirectiveKind::Shutdown:
+ case OpenACCDirectiveKind::Update:
+ case OpenACCDirectiveKind::Invalid:
return false;
case OpenACCDirectiveKind::Parallel:
case OpenACCDirectiveKind::Serial:
@@ -586,6 +592,7 @@ bool doesDirectiveHaveAssociatedStmt(OpenACCDirectiveKind DirKind) {
case OpenACCDirectiveKind::Loop:
case OpenACCDirectiveKind::Data:
case OpenACCDirectiveKind::HostData:
+ case OpenACCDirectiveKind::Atomic:
return true;
}
llvm_unreachable("Unhandled directive->assoc stmt");
@@ -1428,6 +1435,7 @@ Parser::ParseOpenACCDirective() {
SourceLocation DirLoc = getCurToken().getLocation();
OpenACCDirectiveKind DirKind = ParseOpenACCDirectiveKind(*this);
Parser::OpenACCWaitParseInfo WaitInfo;
+ OpenACCAtomicKind AtomicKind = OpenACCAtomicKind::None;
getActions().OpenACC().ActOnConstruct(DirKind, DirLoc);
@@ -1435,7 +1443,7 @@ Parser::ParseOpenACCDirective() {
// specifiers that need to be taken care of. Atomic has an 'atomic-clause'
// that needs to be parsed.
if (DirKind == OpenACCDirectiveKind::Atomic)
- ParseOpenACCAtomicKind(*this);
+ AtomicKind = ParseOpenACCAtomicKind(*this);
// We've successfully parsed the construct/directive name, however a few of
// the constructs have optional parens that contain further details.
@@ -1490,6 +1498,7 @@ Parser::ParseOpenACCDirective() {
T.getCloseLocation(),
/*EndLoc=*/SourceLocation{},
WaitInfo.QueuesLoc,
+ AtomicKind,
WaitInfo.getAllExprs(),
ParseOpenACCClauseList(DirKind)};
@@ -1538,11 +1547,12 @@ StmtResult Parser::ParseOpenACCDirectiveStmt() {
ParseScope ACCScope(this, getOpenACCScopeFlags(DirInfo.DirKind));
AssocStmt = getActions().OpenACC().ActOnAssociatedStmt(
- DirInfo.StartLoc, DirInfo.DirKind, DirInfo.Clauses, ParseStatement());
+ DirInfo.StartLoc, DirInfo.DirKind, DirInfo.AtomicKind, DirInfo.Clauses,
+ ParseStatement());
}
return getActions().OpenACC().ActOnEndStmtDirective(
DirInfo.DirKind, DirInfo.StartLoc, DirInfo.DirLoc, DirInfo.LParenLoc,
- DirInfo.MiscLoc, DirInfo.Exprs, DirInfo.RParenLoc, DirInfo.EndLoc,
- DirInfo.Clauses, AssocStmt);
+ DirInfo.MiscLoc, DirInfo.Exprs, DirInfo.AtomicKind, DirInfo.RParenLoc,
+ DirInfo.EndLoc, DirInfo.Clauses, AssocStmt);
}
diff --git a/clang/lib/Sema/CMakeLists.txt b/clang/lib/Sema/CMakeLists.txt
index 19cf3a2db00fdcd..1a351684d133ebe 100644
--- a/clang/lib/Sema/CMakeLists.txt
+++ b/clang/lib/Sema/CMakeLists.txt
@@ -71,6 +71,7 @@ add_clang_library(clangSema
SemaObjC.cpp
SemaObjCProperty.cpp
SemaOpenACC.cpp
+ SemaOpenACCAtomic.cpp
SemaOpenACCClause.cpp
SemaOpenCL.cpp
SemaOpenMP.cpp
diff --git a/clang/lib/Sema/SemaExceptionSpec.cpp b/clang/lib/Sema/SemaExceptionSpec.cpp
index 7b08a066d3cc1b2..77a1bbcc74e50b1 100644
--- a/clang/lib/Sema/SemaExceptionSpec.cpp
+++ b/clang/lib/Sema/SemaExceptionSpec.cpp
@@ -1425,6 +1425,7 @@ CanThrowResult Sema::canThrow(const Stmt *S) {
case Stmt::OpenACCCombinedConstructClass:
case Stmt::OpenACCDataConstructClass:
case Stmt::OpenACCHostDataConstructClass:
+ case Stmt::OpenACCAtomicConstructClass:
case Stmt::AttributedStmtClass:
case Stmt::BreakStmtClass:
case Stmt::CapturedStmtClass:
diff --git a/clang/lib/Sema/SemaOpenACC.cpp b/clang/lib/Sema/SemaOpenACC.cpp
index f5edc0ed36a9a0a..2d2f8ddf4652beb 100644
--- a/clang/lib/Sema/SemaOpenACC.cpp
+++ b/clang/lib/Sema/SemaOpenACC.cpp
@@ -30,18 +30,23 @@ bool diagnoseConstructAppertainment(SemaOpenACC &S, OpenACCDirectiveKind K,
// Nothing to do here, both invalid and unimplemented don't really need to
// do anything.
break;
- case OpenACCDirectiveKind::ParallelLoop:
- case OpenACCDirectiveKind::SerialLoop:
- case OpenACCDirectiveKind::KernelsLoop:
case OpenACCDirectiveKind::Parallel:
+ case OpenACCDirectiveKind::ParallelLoop:
case OpenACCDirectiveKind::Serial:
+ case OpenACCDirectiveKind::SerialLoop:
case OpenACCDirectiveKind::Kernels:
+ case OpenACCDirectiveKind::KernelsLoop:
case OpenACCDirectiveKind::Loop:
case OpenACCDirectiveKind::Data:
case OpenACCDirectiveKind::EnterData:
case OpenACCDirectiveKind::ExitData:
case OpenACCDirectiveKind::HostData:
case OpenACCDirectiveKind::Wait:
+ case OpenACCDirectiveKind::Update:
+ case OpenACCDirectiveKind::Init:
+ case OpenACCDirectiveKind::Shutdown:
+ case OpenACCDirectiveKind::Cache:
+ case OpenACCDirectiveKind::Atomic:
if (!IsStmt)
return S.Diag(StartLoc, diag::err_acc_construct_appertainment) << K;
break;
@@ -73,6 +78,7 @@ bool PreserveLoopRAIIDepthInAssociatedStmtRAII(OpenACCDirectiveKind DK) {
return false;
case OpenACCDirectiveKind::Data:
case OpenACCDirectiveKind::HostData:
+ case OpenACCDirectiveKind::Atomic:
return true;
case OpenACCDirectiveKind::EnterData:
case OpenACCDirectiveKind::ExitData:
@@ -327,6 +333,7 @@ void SemaOpenACC::ActOnConstruct(OpenACCDirectiveKind K,
case OpenACCDirectiveKind::Shutdown:
case OpenACCDirectiveKind::Set:
case OpenACCDirectiveKind::Update:
+ case OpenACCDirectiveKind::Atomic:
// Nothing to do here, there is no real legalization that needs to happen
// here as these constructs do not take any arguments.
break;
@@ -1518,8 +1525,9 @@ bool SemaOpenACC::ActOnStartStmtDirective(
StmtResult SemaOpenACC::ActOnEndStmtDirective(
OpenACCDirectiveKind K, SourceLocation StartLoc, SourceLocation DirLoc,
SourceLocation LParenLoc, SourceLocation MiscLoc, ArrayRef<Expr *> Exprs,
- SourceLocation RParenLoc, SourceLocation EndLoc,
- ArrayRef<OpenACCClause *> Clauses, StmtResult AssocStmt) {
+ OpenACCAtomicKind AtomicKind, SourceLocation RParenLoc,
+ SourceLocation EndLoc, ArrayRef<OpenACCClause *> Clauses,
+ StmtResult AssocStmt) {
switch (K) {
default:
return StmtEmpty();
@@ -1583,13 +1591,20 @@ StmtResult SemaOpenACC::ActOnEndStmtDirective(
return OpenACCUpdateConstruct::Create(getASTContext(), StartLoc, DirLoc,
EndLoc, Clauses);
}
+ case OpenACCDirectiveKind::Atomic: {
+ assert(Clauses.empty() && "Atomic doesn't allow clauses");
+ return OpenACCAtomicConstruct::Create(
+ getASTContext(), StartLoc, DirLoc, AtomicKind, EndLoc,
+ AssocStmt.isUsable() ? AssocStmt.get() : nullptr);
+ }
}
llvm_unreachable("Unhandled case in directive handling?");
}
StmtResult SemaOpenACC::ActOnAssociatedStmt(
SourceLocation DirectiveLoc, OpenACCDirectiveKind K,
- ArrayRef<const OpenACCClause *> Clauses, StmtResult AssocStmt) {
+ OpenACCAtomicKind AtKind, ArrayRef<const OpenACCClause *> Clauses,
+ StmtResult AssocStmt) {
switch (K) {
default:
llvm_unreachable("Unimplemented associated statement application");
@@ -1601,6 +1616,8 @@ StmtResult SemaOpenACC::ActOnAssociatedStmt(
case OpenACCDirectiveKind::Set:
llvm_unreachable(
"these don't have associated statements, so shouldn't get here");
+ case OpenACCDirectiveKind::Atomic:
+ return CheckAtomicAssociatedStmt(DirectiveLoc, AtKind, AssocStmt);
case OpenACCDirectiveKind::Parallel:
case OpenACCDirectiveKind::Serial:
case OpenACCDirectiveKind::Kernels:
diff --git a/clang/lib/Sema/SemaOpenACCAtomic.cpp b/clang/lib/Sema/SemaOpenACCAtomic.cpp
new file mode 100644
index 000000000000000..68cf338c0711536
--- /dev/null
+++ b/clang/lib/Sema/SemaOpenACCAtomic.cpp
@@ -0,0 +1,736 @@
+//== SemaOpenACCAtomic.cpp - Semantic Analysis for OpenACC Atomic Construct===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+/// \file
+/// This file implements semantic analysis for the OpenACC atomic construct.
+///
+//===----------------------------------------------------------------------===//
+
+#include "clang/AST/ExprCXX.h"
+#include "clang/Basic/DiagnosticSema.h"
+#include "clang/Sema/SemaOpenACC.h"
+
+#include <optional>
+#include <variant>
+
+using namespace clang;
+
+namespace {
+
+class AtomicOperandChecker {
+ SemaOpenACC &SemaRef;
+ OpenACCAtomicKind AtKind;
+ SourceLocation AtomicDirLoc;
+ StmtResult AssocStmt;
+
+ // Do a diagnostic, which sets the correct error, then displays passed note.
+ bool DiagnoseInvalidAtomic(SourceLocation Loc, PartialDiagnostic NoteDiag) {
+ SemaRef.Diag(AtomicDirLoc, diag::err_acc_invalid_atomic)
+ << (AtKind != OpenACCAtomicKind::None) << AtKind;
+ SemaRef.Diag(Loc, NoteDiag);
+ return true;
+ }
+
+ // Create a replacement recovery expr in case we find an error here. This
+ // allows us to ignore this during template instantiation so we only get a
+ // single error.
+ StmtResult getRecoveryExpr() {
+ if (!AssocStmt.isUsable())
+ return AssocStmt;
+
+ if (!SemaRef.getASTContext().getLangOpts().RecoveryAST)
+ return StmtError();
+
+ Expr *E = dyn_cast<Expr>(AssocStmt.get());
+ QualType T = E ? E->getType() : SemaRef.getASTContext().DependentTy;
+
+ return RecoveryExpr::Create(SemaRef.getASTContext(), T,
+ AssocStmt.get()->getBeginLoc(),
+ AssocStmt.get()->getEndLoc(),
+ E ? ArrayRef<Expr *>{E} : ArrayRef<Expr *>{});
+ }
+
+ // OpenACC 3.3 2.12: 'expr' is an expression with scalar type.
+ bool CheckOperandExpr(const Expr *E, PartialDiagnostic PD) {
+ QualType ExprTy = E->getType();
+
+ // Scalar allowed, plus we allow instantiation dependent to support
+ // templates.
+ if (ExprTy->isInstantiationDependentType() || ExprTy->isScalarType())
+ return false;
+
+ return DiagnoseInvalidAtomic(E->getExprLoc(),
+ PD << diag::OACCLValScalar::Scalar << ExprTy);
+ }
+
+ // OpenACC 3.3 2.12: 'x' and 'v' (as applicable) are boht l-value expressoins
+ // with scalar type.
+ bool CheckOperandVariable(const Expr *E, PartialDiagnostic PD) {
+ if (CheckOperandExpr(E, PD))
+ return true;
+
+ if (E->isLValue())
+ return false;
+
+ return DiagnoseInvalidAtomic(E->getExprLoc(),
+ PD << diag::OACCLValScalar::LVal);
+ }
+
+ Expr *RequireExpr(Stmt *Stmt, PartialDiagnostic ExpectedNote) {
+ if (Expr *E = dyn_cast<Expr>(Stmt))
+ return E->IgnoreImpCasts();
+
+ DiagnoseInvalidAtomic(Stmt->getBeginLoc(), ExpectedNote);
+ return nullptr;
+ }
+
+ // A struct to hold the return the inner components of any operands, which
+ // allows for compound checking.
+ struct BinaryOpInfo {
+ const Expr *FoundExpr = nullptr;
+ const Expr *LHS = nullptr;
+ const Expr *RHS = nullptr;
+ BinaryOperatorKind Operator;
+ };
+
+ struct UnaryOpInfo {
+ const Expr *FoundExpr = nullptr;
+ const Expr *SubExpr = nullptr;
+ UnaryOperatorKind Operator;
+
+ bool IsIncrementOp() {
+ return Operator == UO_PostInc || Operator == UO_PreInc;
+ }
+ };
+
+ std::optional<UnaryOpInfo> GetUnaryOperatorInfo(const Expr *E) {
+ // If this is a simple unary operator, just return its details.
+ if (const auto *UO = dyn_cast<UnaryOperator>(E))
+ return UnaryOpInfo{UO, UO->getSubExpr()->IgnoreImpCasts(),
+ UO->getOpcode()};
+
+ // This might be an overloaded operator or a dependent context, so make sure
+ // we can get as many details out of this as we can.
+ if (const auto *OpCall = dyn_cast<CXXOperatorCallExpr>(E)) {
+ UnaryOpInfo Inf;
+ Inf.FoundExpr = OpCall;
+
+ switch (OpCall->getOperator()) {
+ default:
+ return std::nullopt;
+ case OO_PlusPlus:
+ Inf.Operator = OpCall->getNumArgs() == 1 ? UO_PreInc : UO_PostInc;
+ break;
+ case OO_MinusMinus:
+ Inf.Operator = OpCall->getNumArgs() == 1 ? UO_PreDec : UO_PostDec;
+ break;
+ case OO_Amp:
+ Inf.Operator = UO_AddrOf;
+ break;
+ case OO_Star:
+ Inf.Operator = UO_Deref;
+ break;
+ case OO_Plus:
+ Inf.Operator = UO_Plus;
+ break;
+ case OO_Minus:
+ Inf.Operator = UO_Minus;
+ break;
+ case OO_Tilde:
+ Inf.Operator = UO_Not;
+ break;
+ case OO_Exclaim:
+ Inf.Operator = UO_LNot;
+ break;
+ case OO_Coawait:
+ Inf.Operator = UO_Coawait;
+ break;
+ }
+
+ // Some of the above can be both binary and unary operations, so make sure
+ // we get the right one.
+ if (Inf.Operator != UO_PostInc && Inf.Operator != UO_PostDec &&
+ OpCall->getNumArgs() != 1)
+ return std::nullopt;
+
+ Inf.SubExpr = OpCall->getArg(0);
+ return Inf;
+ }
+ return std::nullopt;
+ }
+
+ // Get a normalized version of a binary operator.
+ std::optional<BinaryOpInfo> GetBinaryOperatorInfo(const Expr *E) {
+ if (const auto *BO = dyn_cast<BinaryOperator>(E))
+ return BinaryOpInfo{BO, BO->getLHS()->IgnoreImpCasts(),
+ BO->getRHS()->IgnoreImpCasts(), BO->getOpcode()};
+
+ // In case this is an operator-call, which allows us to support overloaded
+ // operators and dependent expression.
+ if (const auto *OpCall = dyn_cast<CXXOperatorCallExpr>(E)) {
+ BinaryOpInfo Inf;
+ Inf.FoundExpr = OpCall;
+
+ switch (OpCall->getOperator()) {
+ default:
+ return std::nullopt;
+ case OO_Plus:
+ Inf.Operator = BO_Add;
+ break;
+ case OO_Minus:
+ Inf.Operator = BO_Sub;
+ break;
+ case OO_Star:
+ Inf.Operator = BO_Mul;
+ break;
+ case OO_Slash:
+ Inf.Operator = BO_Div;
+ break;
+ case OO_Percent:
+ Inf.Operator = BO_Rem;
+ break;
+ case OO_Caret:
+ Inf.Operator = BO_Xor;
+ break;
+ case OO_Amp:
+ Inf.Operator = BO_And;
+ break;
+ case OO_Pipe:
+ Inf.Operator = BO_Or;
+ break;
+ case OO_Equal:
+ Inf.Operator = BO_Assign;
+ break;
+ case OO_Spaceship:
+ Inf.Operator = BO_Cmp;
+ break;
+ case OO_Less:
+ Inf.Operator = BO_LT;
+ break;
+ case OO_Greater:
+ Inf.Operator = BO_GT;
+ break;
+ case OO_PlusEqual:
+ Inf.Operator = BO_AddAssign;
+ break;
+ case OO_MinusEqual:
+ Inf.Operator = BO_SubAssign;
+ break;
+ case OO_StarEqual:
+ Inf.Operator = BO_MulAssign;
+ break;
+ case OO_SlashEqual:
+ Inf.Operator = BO_DivAssign;
+ break;
+ case OO_PercentEqual:
+ Inf.Operator = BO_RemAssign;
+ break;
+ case OO_CaretEqual:
+ Inf.Operator = BO_XorAssign;
+ break;
+ case OO_AmpEqual:
+ Inf.Operator = BO_AndAssign;
+ break;
+ case OO_PipeEqual:
+ Inf.Operator = BO_OrAssign;
+ break;
+ case OO_LessLess:
+ Inf.Operator = BO_Shl;
+ break;
+ case OO_GreaterGreater:
+ Inf.Operator = BO_Shr;
+ break;
+ case OO_LessLessEqual:
+ Inf.Operator = BO_ShlAssign;
+ break;
+ case OO_GreaterGreaterEqual:
+ Inf.Operator = BO_ShrAssign;
+ break;
+ case OO_EqualEqual:
+ Inf.Operator = BO_EQ;
+ break;
+ case OO_ExclaimEqual:
+ Inf.Operator = BO_NE;
+ break;
+ case OO_LessEqual:
+ Inf.Operator = BO_LE;
+ break;
+ case OO_GreaterEqual:
+ Inf.Operator = BO_GE;
+ break;
+ case OO_AmpAmp:
+ Inf.Operator = BO_LAnd;
+ break;
+ case OO_PipePipe:
+ Inf.Operator = BO_LOr;
+ break;
+ case OO_Comma:
+ Inf.Operator = BO_Comma;
+ break;
+ case OO_ArrowStar:
+ Inf.Operator = BO_PtrMemI;
+ break;
+ }
+
+ // This isn't a binary operator unless there are two arguments.
+ if (OpCall->getNumArgs() != 2)
+ return std::nullopt;
+
+ // Callee is the call-operator, so we only need to extract the two
+ // arguments here.
+ Inf.LHS = OpCall->getArg(0)->IgnoreImpCasts();
+ Inf.RHS = OpCall->getArg(1)->IgnoreImpCasts();
+ return Inf;
+ }
+
+ return std::nullopt;
+ }
+
+ // Checks a required assignment operation, but don't check the LHS or RHS,
+ // callers have to do that here.
+ std::optional<BinaryOpInfo> CheckAssignment(const Expr *E) {
+ std::optional<BinaryOpInfo> Inf = GetBinaryOperatorInfo(E);
+
+ if (!Inf) {
+ DiagnoseInvalidAtomic(E->getExprLoc(),
+ SemaRef.PDiag(diag::note_acc_atomic_expr_must_be)
+ << diag::OACCAtomicExpr::Assign);
+ return std::nullopt;
+ }
+
+ if (Inf->Operator != BO_Assign) {
+ DiagnoseInvalidAtomic(Inf->FoundExpr->getExprLoc(),
+ SemaRef.PDiag(diag::note_acc_atomic_expr_must_be)
+ << diag::OACCAtomicExpr::Assign);
+ return std::nullopt;
+ }
+
+ // Assignment always requires an lvalue/scalar on the LHS.
+ if (CheckOperandVariable(
+ Inf->LHS, SemaRef.PDiag(diag::note_acc_atomic_operand_lvalue_scalar)
+ << /*left=*/0 << diag::OACCAtomicOpKind::Assign))
+ return std::nullopt;
+
+ return Inf;
+ }
+
+ struct IDACInfo {
+ bool Failed = false;
+ enum ExprKindTy {
+ Invalid,
+ // increment/decrement ops.
+ Unary,
+ // v = x
+ SimpleAssign,
+ // x = expr
+ ExprAssign,
+ // x binop= expr
+ CompoundAssign,
+ // x = x binop expr
+ // x = expr binop x
+ AssignBinOp
+ } ExprKind;
+
+ // The variable referred to as 'x' in all of the grammar, such that it is
+ // needed in compound statement checking of capture to check between the two
+ // expressions.
+ const Expr *X_Var = nullptr;
+
+ static IDACInfo Fail() { return IDACInfo{true, Invalid, nullptr}; };
+ };
+
+ // Helper for CheckIncDecAssignCompoundAssign, does checks for inc/dec.
+ IDACInfo CheckIncDec(UnaryOpInfo Inf) {
+
+ if (!UnaryOperator::isIncrementDecrementOp(Inf.Operator)) {
+ DiagnoseInvalidAtomic(
+ Inf.FoundExpr->getExprLoc(),
+ SemaRef.PDiag(diag::note_acc_atomic_unsupported_unary_operator));
+ return IDACInfo::Fail();
+ }
+ bool Failed = CheckOperandVariable(
+ Inf.SubExpr,
+ SemaRef.PDiag(diag::note_acc_atomic_operand_lvalue_scalar)
+ << /*none=*/2
+ << (Inf.IsIncrementOp() ? diag::OACCAtomicOpKind::Inc
+ : diag::OACCAtomicOpKind::Dec));
+ // For increment/decrements, the subexpr is the 'x' (x++, ++x, etc).
+ return IDACInfo{Failed, IDACInfo::Unary, Inf.SubExpr};
+ }
+
+ enum class SimpleAssignKind { None, Var, Expr };
+
+ // Check an assignment, and ensure the RHS is either x binop expr or expr
+ // binop x.
+ // If AllowSimpleAssign, also allows v = x;
+ IDACInfo CheckAssignmentWithBinOpOnRHS(BinaryOpInfo AssignInf,
+ SimpleAssignKind SAK) {
+ PartialDiagnostic PD =
+ SemaRef.PDiag(diag::note_acc_atomic_operand_lvalue_scalar)
+ << /*left=*/0 << diag::OACCAtomicOpKind::Assign;
+ if (CheckOperandVariable(AssignInf.LHS, PD))
+ return IDACInfo::Fail();
+
+ std::optional<BinaryOpInfo> BinInf = GetBinaryOperatorInfo(AssignInf.RHS);
+
+ if (!BinInf) {
+
+ // Capture in a compound statement allows v = x assignment. So make sure
+ // we permit that here.
+ if (SAK != SimpleAssignKind::None) {
+ PartialDiagnostic PD =
+ SemaRef.PDiag(diag::note_acc_atomic_operand_lvalue_scalar)
+ << /*right=*/1 << diag::OACCAtomicOpKind::Assign;
+ if (SAK == SimpleAssignKind::Var) {
+ // In the var version, everywhere we allow v = x;, X is the RHS.
+ return IDACInfo{CheckOperandVariable(AssignInf.RHS, PD),
+ IDACInfo::SimpleAssign, AssignInf.RHS};
+ }
+ assert(SAK == SimpleAssignKind::Expr);
+ // In the expression version, supported by v=x; x = expr;, we need to
+ // set to the LHS here.
+ return IDACInfo{CheckOperandExpr(AssignInf.RHS, PD),
+ IDACInfo::ExprAssign, AssignInf.LHS};
+ }
+
+ DiagnoseInvalidAtomic(
+ AssignInf.RHS->getExprLoc(),
+ SemaRef.PDiag(diag::note_acc_atomic_expected_binop));
+
+ return IDACInfo::Fail();
+ }
+ switch (BinInf->Operator) {
+ default:
+ DiagnoseInvalidAtomic(
+ BinInf->FoundExpr->getExprLoc(),
+ SemaRef.PDiag(diag::note_acc_atomic_unsupported_binary_operator));
+ return IDACInfo::Fail();
+ // binop is one of +, *, -, /, &, ^, |, <<, or >>
+ case BO_Add:
+ case BO_Mul:
+ case BO_Sub:
+ case BO_Div:
+ case BO_And:
+ case BO_Xor:
+ case BO_Or:
+ case BO_Shl:
+ case BO_Shr:
+ // Handle these outside of the switch.
+ break;
+ }
+
+ llvm::FoldingSetNodeID LHS_ID, InnerLHS_ID, InnerRHS_ID;
+ AssignInf.LHS->Profile(LHS_ID, SemaRef.getASTContext(),
+ /*Canonical=*/true);
+ BinInf->LHS->Profile(InnerLHS_ID, SemaRef.getASTContext(),
+ /*Canonical=*/true);
+
+ // This is X = X binop expr;
+ // Check the RHS is an expression.
+ if (LHS_ID == InnerLHS_ID)
+ return IDACInfo{
+ CheckOperandExpr(
+ BinInf->RHS,
+ SemaRef.PDiag(diag::note_acc_atomic_operand_lvalue_scalar
+ << /*right=*/1
+ << diag::OACCAtomicOpKind::CompoundAssign)),
+ IDACInfo::AssignBinOp, AssignInf.LHS};
+
+ BinInf->RHS->Profile(InnerRHS_ID, SemaRef.getASTContext(),
+ /*Canonical=*/true);
+ // This is X = expr binop X;
+ // Check the LHS is an expression
+ if (LHS_ID == InnerRHS_ID)
+ return IDACInfo{
+ CheckOperandExpr(
+ BinInf->LHS,
+ SemaRef.PDiag(diag::note_acc_atomic_operand_lvalue_scalar)
+ << /*left=*/0 << diag::OACCAtomicOpKind::CompoundAssign),
+ IDACInfo::AssignBinOp, AssignInf.LHS};
+
+ // If nothing matches, error out.
+ DiagnoseInvalidAtomic(BinInf->FoundExpr->getExprLoc(),
+ SemaRef.PDiag(diag::note_acc_atomic_mismatch_operand)
+ << const_cast<Expr *>(AssignInf.LHS)
+ << const_cast<Expr *>(BinInf->LHS)
+ << const_cast<Expr *>(BinInf->RHS));
+ return IDACInfo::Fail();
+ }
+
+ // Ensures that the expression is an increment/decrement, an assignment, or a
+ // compound assignment. If its an assignment, allows the x binop expr/x binop
+ // expr syntax. If it is a compound-assignment, allows any expr on the RHS.
+ IDACInfo CheckIncDecAssignCompoundAssign(const Expr *E,
+ SimpleAssignKind SAK) {
+ std::optional<UnaryOpInfo> UInf = GetUnaryOperatorInfo(E);
+
+ // If this is a unary operator, only increment/decrement are allowed, so get
+ // unary operator, then check everything we can.
+ if (UInf)
+ return CheckIncDec(*UInf);
+
+ std::optional<BinaryOpInfo> BinInf = GetBinaryOperatorInfo(E);
+
+ // Unary or binary operator were the only choices, so error here.
+ if (!BinInf) {
+ DiagnoseInvalidAtomic(E->getExprLoc(),
+ SemaRef.PDiag(diag::note_acc_atomic_expr_must_be)
+ << diag::OACCAtomicExpr::UnaryCompAssign);
+ return IDACInfo::Fail();
+ }
+
+ switch (BinInf->Operator) {
+ default:
+ DiagnoseInvalidAtomic(
+ BinInf->FoundExpr->getExprLoc(),
+ SemaRef.PDiag(
+ diag::note_acc_atomic_unsupported_compound_binary_operator));
+ return IDACInfo::Fail();
+ case BO_Assign:
+ return CheckAssignmentWithBinOpOnRHS(*BinInf, SAK);
+ case BO_AddAssign:
+ case BO_MulAssign:
+ case BO_SubAssign:
+ case BO_DivAssign:
+ case BO_AndAssign:
+ case BO_XorAssign:
+ case BO_OrAssign:
+ case BO_ShlAssign:
+ case BO_ShrAssign: {
+ PartialDiagnostic LPD =
+ SemaRef.PDiag(diag::note_acc_atomic_operand_lvalue_scalar)
+ << /*left=*/0 << diag::OACCAtomicOpKind::CompoundAssign;
+ PartialDiagnostic RPD =
+ SemaRef.PDiag(diag::note_acc_atomic_operand_lvalue_scalar)
+ << /*right=*/1 << diag::OACCAtomicOpKind::CompoundAssign;
+ // nothing to do other than check the variable expressions.
+ // success or failure
+ bool Failed = CheckOperandVariable(BinInf->LHS, LPD) ||
+ CheckOperandExpr(BinInf->RHS, RPD);
+
+ return IDACInfo{Failed, IDACInfo::CompoundAssign, BinInf->LHS};
+ }
+ }
+ llvm_unreachable("all binary operator kinds should be checked above");
+ }
+
+ StmtResult CheckRead() {
+ Expr *AssocExpr = RequireExpr(
+ AssocStmt.get(), SemaRef.PDiag(diag::note_acc_atomic_expr_must_be)
+ << diag::OACCAtomicExpr::Assign);
+
+ if (!AssocExpr)
+ return getRecoveryExpr();
+
+ std::optional<BinaryOpInfo> AssignRes = CheckAssignment(AssocExpr);
+ if (!AssignRes)
+ return getRecoveryExpr();
+
+ PartialDiagnostic PD =
+ SemaRef.PDiag(diag::note_acc_atomic_operand_lvalue_scalar)
+ << /*right=*/1 << diag::OACCAtomicOpKind::Assign;
+
+ // Finally, check the RHS.
+ if (CheckOperandVariable(AssignRes->RHS, PD))
+ return getRecoveryExpr();
+
+ return AssocStmt;
+ }
+
+ StmtResult CheckWrite() {
+ Expr *AssocExpr = RequireExpr(
+ AssocStmt.get(), SemaRef.PDiag(diag::note_acc_atomic_expr_must_be)
+ << diag::OACCAtomicExpr::Assign);
+
+ if (!AssocExpr)
+ return getRecoveryExpr();
+
+ std::optional<BinaryOpInfo> AssignRes = CheckAssignment(AssocExpr);
+ if (!AssignRes)
+ return getRecoveryExpr();
+
+ PartialDiagnostic PD =
+ SemaRef.PDiag(diag::note_acc_atomic_operand_lvalue_scalar)
+ << /*right=*/1 << diag::OACCAtomicOpKind::Assign;
+
+ // Finally, check the RHS.
+ if (CheckOperandExpr(AssignRes->RHS, PD))
+ return getRecoveryExpr();
+
+ return AssocStmt;
+ }
+
+ StmtResult CheckUpdate() {
+ Expr *AssocExpr = RequireExpr(
+ AssocStmt.get(), SemaRef.PDiag(diag::note_acc_atomic_expr_must_be)
+ << diag::OACCAtomicExpr::UnaryCompAssign);
+
+ if (!AssocExpr ||
+ CheckIncDecAssignCompoundAssign(AssocExpr, SimpleAssignKind::None)
+ .Failed)
+ return getRecoveryExpr();
+
+ return AssocStmt;
+ }
+
+ bool CheckVarRefsSame(IDACInfo::ExprKindTy FirstKind, const Expr *FirstX,
+ IDACInfo::ExprKindTy SecondKind, const Expr *SecondX) {
+ llvm::FoldingSetNodeID First_ID, Second_ID;
+ FirstX->Profile(First_ID, SemaRef.getASTContext(), /*Canonical=*/true);
+ SecondX->Profile(Second_ID, SemaRef.getASTContext(), /*Canonical=*/true);
+
+ if (First_ID == Second_ID)
+ return false;
+
+ PartialDiagnostic PD =
+ SemaRef.PDiag(diag::note_acc_atomic_mismatch_compound_operand)
+ << FirstKind << const_cast<Expr *>(FirstX) << SecondKind
+ << const_cast<Expr *>(SecondX);
+
+ return DiagnoseInvalidAtomic(SecondX->getExprLoc(), PD);
+ }
+
+ StmtResult CheckCapture() {
+ if (const auto *CmpdStmt = dyn_cast<CompoundStmt>(AssocStmt.get())) {
+ auto *const *BodyItr = CmpdStmt->body().begin();
+ PartialDiagnostic PD = SemaRef.PDiag(diag::note_acc_atomic_expr_must_be)
+ << diag::OACCAtomicExpr::UnaryCompAssign;
+ // If we don't have at least 1 statement, error.
+ if (BodyItr == CmpdStmt->body().end()) {
+ DiagnoseInvalidAtomic(CmpdStmt->getBeginLoc(), PD);
+ return getRecoveryExpr();
+ }
+
+ // First Expr can be inc/dec, assign, or compound assign.
+ Expr *FirstExpr = RequireExpr(*BodyItr, PD);
+ if (!FirstExpr)
+ return getRecoveryExpr();
+
+ IDACInfo FirstExprResults =
+ CheckIncDecAssignCompoundAssign(FirstExpr, SimpleAssignKind::Var);
+ if (FirstExprResults.Failed)
+ return getRecoveryExpr();
+
+ ++BodyItr;
+
+ // If we don't have second statement, error.
+ if (BodyItr == CmpdStmt->body().end()) {
+ DiagnoseInvalidAtomic(CmpdStmt->getEndLoc(), PD);
+ return getRecoveryExpr();
+ }
+
+ Expr *SecondExpr = RequireExpr(*BodyItr, PD);
+ if (!SecondExpr)
+ return getRecoveryExpr();
+
+ assert(FirstExprResults.ExprKind != IDACInfo::Invalid);
+
+ switch (FirstExprResults.ExprKind) {
+ case IDACInfo::Invalid:
+ case IDACInfo::ExprAssign:
+ llvm_unreachable("Should have error'ed out by now");
+ case IDACInfo::Unary:
+ case IDACInfo::CompoundAssign:
+ case IDACInfo::AssignBinOp: {
+ // Everything but simple-assign can only be followed by a simple
+ // assignment.
+ std::optional<BinaryOpInfo> AssignRes = CheckAssignment(SecondExpr);
+ if (!AssignRes)
+ return getRecoveryExpr();
+
+ PartialDiagnostic PD =
+ SemaRef.PDiag(diag::note_acc_atomic_operand_lvalue_scalar)
+ << /*right=*/1 << diag::OACCAtomicOpKind::Assign;
+
+ if (CheckOperandVariable(AssignRes->RHS, PD))
+ return getRecoveryExpr();
+
+ if (CheckVarRefsSame(FirstExprResults.ExprKind, FirstExprResults.X_Var,
+ IDACInfo::SimpleAssign, AssignRes->RHS))
+ return getRecoveryExpr();
+ break;
+ }
+ case IDACInfo::SimpleAssign: {
+ // If the first was v = x, anything but simple expression is allowed.
+ IDACInfo SecondExprResults =
+ CheckIncDecAssignCompoundAssign(SecondExpr, SimpleAssignKind::Expr);
+ if (SecondExprResults.Failed)
+ return getRecoveryExpr();
+
+ if (CheckVarRefsSame(FirstExprResults.ExprKind, FirstExprResults.X_Var,
+ SecondExprResults.ExprKind,
+ SecondExprResults.X_Var))
+ return getRecoveryExpr();
+ break;
+ }
+ }
+ ++BodyItr;
+ if (BodyItr != CmpdStmt->body().end()) {
+ DiagnoseInvalidAtomic(
+ (*BodyItr)->getBeginLoc(),
+ SemaRef.PDiag(diag::note_acc_atomic_too_many_stmts));
+ return getRecoveryExpr();
+ }
+ } else {
+ // This check doesn't need to happen if it is a compound stmt.
+ Expr *AssocExpr = RequireExpr(
+ AssocStmt.get(), SemaRef.PDiag(diag::note_acc_atomic_expr_must_be)
+ << diag::OACCAtomicExpr::Assign);
+ if (!AssocExpr)
+ return getRecoveryExpr();
+
+ // First, we require an assignment.
+ std::optional<BinaryOpInfo> AssignRes = CheckAssignment(AssocExpr);
+
+ if (!AssignRes)
+ return getRecoveryExpr();
+
+ if (CheckIncDecAssignCompoundAssign(AssignRes->RHS,
+ SimpleAssignKind::None)
+ .Failed)
+ return getRecoveryExpr();
+ }
+
+ return AssocStmt;
+ }
+
+public:
+ AtomicOperandChecker(SemaOpenACC &S, OpenACCAtomicKind AtKind,
+ SourceLocation DirLoc, StmtResult AssocStmt)
+ : SemaRef(S), AtKind(AtKind), AtomicDirLoc(DirLoc), AssocStmt(AssocStmt) {
+ }
+
+ StmtResult Check() {
+
+ switch (AtKind) {
+ case OpenACCAtomicKind::Read:
+ return CheckRead();
+ case OpenACCAtomicKind::Write:
+ return CheckWrite();
+ case OpenACCAtomicKind::None:
+ case OpenACCAtomicKind::Update:
+ return CheckUpdate();
+ case OpenACCAtomicKind::Capture:
+ return CheckCapture();
+ }
+ llvm_unreachable("Unhandled atomic kind?");
+ }
+};
+} // namespace
+
+StmtResult SemaOpenACC::CheckAtomicAssociatedStmt(SourceLocation AtomicDirLoc,
+ OpenACCAtomicKind AtKind,
+ StmtResult AssocStmt) {
+ if (!AssocStmt.isUsable())
+ return AssocStmt;
+
+ if (isa<RecoveryExpr>(AssocStmt.get()))
+ return AssocStmt;
+
+ AtomicOperandChecker Checker{*this, AtKind, AtomicDirLoc, AssocStmt};
+ return Checker.Check();
+}
diff --git a/clang/lib/Sema/SemaOpenACCClause.cpp b/clang/lib/Sema/SemaOpenACCClause.cpp
index 000934225402ab5..1e74f126c31cee7 100644
--- a/clang/lib/Sema/SemaOpenACCClause.cpp
+++ b/clang/lib/Sema/SemaOpenACCClause.cpp
@@ -589,7 +589,6 @@ bool checkValidAfterDeviceType(
// construct has been implemented.
bool isDirectiveKindImplemented(OpenACCDirectiveKind DK) {
return DK != OpenACCDirectiveKind::Declare &&
- DK != OpenACCDirectiveKind::Atomic &&
DK != OpenACCDirectiveKind::Routine;
}
diff --git a/clang/lib/Sema/TreeTransform.h b/clang/lib/Sema/TreeTransform.h
index 808b56448e1ea6e..60100d77c22a734 100644
--- a/clang/lib/Sema/TreeTransform.h
+++ b/clang/lib/Sema/TreeTransform.h
@@ -4211,6 +4211,17 @@ class TreeTransform {
Exprs, RParenLoc, EndLoc, Clauses, {});
}
+ StmtResult RebuildOpenACCAtomicConstruct(SourceLocation BeginLoc,
+ SourceLocation DirLoc,
+ OpenACCAtomicKind AtKind,
+ SourceLocation EndLoc,
+ StmtResult AssociatedStmt) {
+ return getSema().OpenACC().ActOnEndStmtDirective(
+ OpenACCDirectiveKind::Atomic, BeginLoc, DirLoc, SourceLocation{},
+ SourceLocation{}, {}, AtKind, SourceLocation{}, EndLoc, {},
+ AssociatedStmt);
+ }
+
ExprResult RebuildOpenACCAsteriskSizeExpr(SourceLocation AsteriskLoc) {
return getSema().OpenACC().ActOnOpenACCAsteriskSizeExpr(AsteriskLoc);
}
@@ -12613,6 +12624,29 @@ TreeTransform<Derived>::TransformOpenACCWaitConstruct(OpenACCWaitConstruct *C) {
QueueIdExprs, C->getRParenLoc(), C->getEndLoc(), TransformedClauses);
}
+template <typename Derived>
+StmtResult TreeTransform<Derived>::TransformOpenACCAtomicConstruct(
+ OpenACCAtomicConstruct *C) {
+ getSema().OpenACC().ActOnConstruct(C->getDirectiveKind(), C->getBeginLoc());
+
+ if (getSema().OpenACC().ActOnStartStmtDirective(C->getDirectiveKind(),
+ C->getBeginLoc(), {}))
+ return StmtError();
+
+ // Transform Associated Stmt.
+ SemaOpenACC::AssociatedStmtRAII AssocStmtRAII(
+ getSema().OpenACC(), C->getDirectiveKind(), C->getDirectiveLoc(), {}, {});
+
+ StmtResult AssocStmt = getDerived().TransformStmt(C->getAssociatedStmt());
+ AssocStmt = getSema().OpenACC().ActOnAssociatedStmt(
+ C->getBeginLoc(), C->getDirectiveKind(), C->getAtomicKind(), {},
+ AssocStmt);
+
+ return getDerived().RebuildOpenACCAtomicConstruct(
+ C->getBeginLoc(), C->getDirectiveLoc(), C->getAtomicKind(),
+ C->getEndLoc(), AssocStmt);
+}
+
template <typename Derived>
ExprResult TreeTransform<Derived>::TransformOpenACCAsteriskSizeExpr(
OpenACCAsteriskSizeExpr *E) {
diff --git a/clang/lib/Serialization/ASTReaderStmt.cpp b/clang/lib/Serialization/ASTReaderStmt.cpp
index b15eca87993a5b3..dc953ddeee85c72 100644
--- a/clang/lib/Serialization/ASTReaderStmt.cpp
+++ b/clang/lib/Serialization/ASTReaderStmt.cpp
@@ -2923,6 +2923,15 @@ void ASTStmtReader::VisitOpenACCWaitConstruct(OpenACCWaitConstruct *S) {
}
}
+void ASTStmtReader::VisitOpenACCAtomicConstruct(OpenACCAtomicConstruct *S) {
+ VisitStmt(S);
+ S->Kind = Record.readEnum<OpenACCDirectiveKind>();
+ S->Range = Record.readSourceRange();
+ S->DirectiveLoc = Record.readSourceLocation();
+ S->AtomicKind = Record.readEnum<OpenACCAtomicKind>();
+ S->setAssociatedStmt(Record.readSubStmt());
+}
+
//===----------------------------------------------------------------------===//
// HLSL Constructs/Directives.
//===----------------------------------------------------------------------===//
@@ -4454,6 +4463,10 @@ Stmt *ASTReader::ReadStmtFromStream(ModuleFile &F) {
S = OpenACCUpdateConstruct::CreateEmpty(Context, NumClauses);
break;
}
+ case STMT_OPENACC_ATOMIC_CONSTRUCT: {
+ S = OpenACCAtomicConstruct::CreateEmpty(Context);
+ break;
+ }
case EXPR_REQUIRES: {
unsigned numLocalParameters = Record[ASTStmtReader::NumExprFields];
unsigned numRequirement = Record[ASTStmtReader::NumExprFields + 1];
diff --git a/clang/lib/Serialization/ASTWriterStmt.cpp b/clang/lib/Serialization/ASTWriterStmt.cpp
index e6701c586e022e7..e5caf3debc02329 100644
--- a/clang/lib/Serialization/ASTWriterStmt.cpp
+++ b/clang/lib/Serialization/ASTWriterStmt.cpp
@@ -3007,6 +3007,17 @@ void ASTStmtWriter::VisitOpenACCWaitConstruct(OpenACCWaitConstruct *S) {
Code = serialization::STMT_OPENACC_WAIT_CONSTRUCT;
}
+void ASTStmtWriter::VisitOpenACCAtomicConstruct(OpenACCAtomicConstruct *S) {
+ VisitStmt(S);
+ Record.writeEnum(S->Kind);
+ Record.AddSourceRange(S->Range);
+ Record.AddSourceLocation(S->DirectiveLoc);
+ Record.writeEnum(S->getAtomicKind());
+ Record.AddStmt(S->getAssociatedStmt());
+
+ Code = serialization::STMT_OPENACC_ATOMIC_CONSTRUCT;
+}
+
//===----------------------------------------------------------------------===//
// HLSL Constructs/Directives.
//===----------------------------------------------------------------------===//
diff --git a/clang/lib/StaticAnalyzer/Core/ExprEngine.cpp b/clang/lib/StaticAnalyzer/Core/ExprEngine.cpp
index 2b1872f8386aad1..9545ce5f2569664 100644
--- a/clang/lib/StaticAnalyzer/Core/ExprEngine.cpp
+++ b/clang/lib/StaticAnalyzer/Core/ExprEngine.cpp
@@ -1836,6 +1836,7 @@ void ExprEngine::Visit(const Stmt *S, ExplodedNode *Pred,
case Stmt::OpenACCShutdownConstructClass:
case Stmt::OpenACCSetConstructClass:
case Stmt::OpenACCUpdateConstructClass:
+ case Stmt::OpenACCAtomicConstructClass:
case Stmt::OMPUnrollDirectiveClass:
case Stmt::OMPMetaDirectiveClass:
case Stmt::HLSLOutArgExprClass: {
diff --git a/clang/test/AST/ast-print-openacc-atomic-construct.cpp b/clang/test/AST/ast-print-openacc-atomic-construct.cpp
new file mode 100644
index 000000000000000..572f2ea4842d88b
--- /dev/null
+++ b/clang/test/AST/ast-print-openacc-atomic-construct.cpp
@@ -0,0 +1,33 @@
+// RUN: %clang_cc1 -fopenacc -ast-print %s -o - | FileCheck %s
+
+void foo(int v, int x) {
+// CHECK: #pragma acc atomic read
+// CHECK-NEXT: v = x;
+#pragma acc atomic read
+ v = x;
+// CHECK-NEXT: pragma acc atomic write
+// CHECK-NEXT: v = x + 1;
+#pragma acc atomic write
+ v = x + 1;
+// CHECK-NEXT: pragma acc atomic update
+// CHECK-NEXT: x++;
+#pragma acc atomic update
+ x++;
+// CHECK-NEXT: pragma acc atomic
+// CHECK-NEXT: x--;
+#pragma acc atomic
+ x--;
+// CHECK-NEXT: pragma acc atomic capture
+// CHECK-NEXT: v = x++;
+#pragma acc atomic capture
+ v = x++;
+
+// CHECK-NEXT: #pragma acc atomic capture
+// CHECK-NEXT: {
+// CHECK-NEXT: x--;
+// CHECK-NEXT: v = x;
+// CHECK-NEXT: }
+#pragma acc atomic capture
+ { x--; v = x; }
+
+}
diff --git a/clang/test/ParserOpenACC/parse-constructs.c b/clang/test/ParserOpenACC/parse-constructs.c
index 886a912713c58be..f0698495a3cc2fc 100644
--- a/clang/test/ParserOpenACC/parse-constructs.c
+++ b/clang/test/ParserOpenACC/parse-constructs.c
@@ -109,30 +109,23 @@ void func() {
for(int i = 0; i < 6;++i){}
int i = 0, j = 0, k = 0;
- // expected-warning at +1{{OpenACC construct 'atomic' not yet implemented, pragma ignored}}
#pragma acc atomic
- i = j;
- // expected-error at +2{{invalid OpenACC clause 'garbage'}}
- // expected-warning at +1{{OpenACC construct 'atomic' not yet implemented, pragma ignored}}
+ i = i + 1;
+ // expected-error at +1{{invalid OpenACC clause 'garbage'}}
#pragma acc atomic garbage
- i = j;
- // expected-error at +2{{invalid OpenACC clause 'garbage'}}
- // expected-warning at +1{{OpenACC construct 'atomic' not yet implemented, pragma ignored}}
+ i = i + 1;
+ // expected-error at +1{{invalid OpenACC clause 'garbage'}}
#pragma acc atomic garbage clause list
- i = j;
- // expected-warning at +1{{OpenACC construct 'atomic' not yet implemented, pragma ignored}}
+ i = i + 1;
#pragma acc atomic read
i = j;
- // expected-error at +2{{invalid OpenACC clause 'clause'}}
- // expected-warning at +1{{OpenACC construct 'atomic' not yet implemented, pragma ignored}}
+ // expected-error at +1{{invalid OpenACC clause 'clause'}}
#pragma acc atomic write clause list
i = i + j;
- // expected-error at +2{{invalid OpenACC clause 'clause'}}
- // expected-warning at +1{{OpenACC construct 'atomic' not yet implemented, pragma ignored}}
+ // expected-error at +1{{invalid OpenACC clause 'clause'}}
#pragma acc atomic update clause list
i++;
- // expected-error at +2{{invalid OpenACC clause 'clause'}}
- // expected-warning at +1{{OpenACC construct 'atomic' not yet implemented, pragma ignored}}
+ // expected-error at +1{{invalid OpenACC clause 'clause'}}
#pragma acc atomic capture clause list
i = j++;
diff --git a/clang/test/SemaOpenACC/atomic-construct-ast.cpp b/clang/test/SemaOpenACC/atomic-construct-ast.cpp
new file mode 100644
index 000000000000000..6579b87941e5a63
--- /dev/null
+++ b/clang/test/SemaOpenACC/atomic-construct-ast.cpp
@@ -0,0 +1,170 @@
+// RUN: %clang_cc1 %s -fopenacc -ast-dump | FileCheck %s
+
+// Test this with PCH.
+// RUN: %clang_cc1 %s -fopenacc -emit-pch -o %t %s
+// RUN: %clang_cc1 %s -fopenacc -include-pch %t -ast-dump-all | FileCheck %s
+
+#ifndef PCH_HELPER
+#define PCH_HELPER
+
+void foo(int v, int x) {
+ // CHECK: FunctionDecl{{.*}} foo 'void (int, int)'
+ // CHECK-NEXT: ParmVarDecl
+ // CHECK-NEXT: ParmVarDecl
+ // CHECK-NEXT: CompoundStmt
+
+// CHECK-NEXT: OpenACCAtomicConstruct{{.*}} atomic read
+// CHECK-NEXT: BinaryOperator{{.*}} 'int' lvalue '='
+// CHECK-NEXT: DeclRefExpr{{.*}}'v' 'int'
+// CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <LValueToRValue>
+// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'int'
+#pragma acc atomic read
+ v = x;
+
+// CHECK-NEXT: OpenACCAtomicConstruct{{.*}} atomic write
+// CHECK-NEXT: BinaryOperator{{.*}} 'int' lvalue '='
+// CHECK-NEXT: DeclRefExpr{{.*}}'v' 'int'
+// CHECK-NEXT: BinaryOperator{{.*}}'int' '+'
+// CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <LValueToRValue>
+// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'int'
+// CHECK-NEXT: IntegerLiteral{{.*}} 'int' 1
+#pragma acc atomic write
+ v = x + 1;
+
+// CHECK-NEXT: OpenACCAtomicConstruct{{.*}} atomic update
+// CHECK-NEXT: UnaryOperator{{.*}} 'int' postfix '++'
+// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'int'
+#pragma acc atomic update
+ x++;
+// CHECK-NEXT: OpenACCAtomicConstruct{{.*}} atomic <none>
+// CHECK-NEXT: UnaryOperator{{.*}} 'int' postfix '--'
+// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'int'
+#pragma acc atomic
+ x--;
+// CHECK-NEXT: OpenACCAtomicConstruct{{.*}} atomic capture
+// CHECK-NEXT: BinaryOperator{{.*}} 'int' lvalue '='
+// CHECK-NEXT: DeclRefExpr{{.*}}'v' 'int'
+// CHECK-NEXT: UnaryOperator{{.*}} 'int' postfix '++'
+// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'int'
+#pragma acc atomic capture
+ v = x++;
+
+// CHECK-NEXT: OpenACCAtomicConstruct{{.*}} atomic capture
+// CHECK-NEXT: CompoundStmt
+// CHECK-NEXT: UnaryOperator{{.*}} 'int' postfix '--'
+// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'int'
+// CHECK-NEXT: BinaryOperator{{.*}} 'int' lvalue '='
+// CHECK-NEXT: DeclRefExpr{{.*}}'v' 'int'
+// CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <LValueToRValue>
+// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'int'
+#pragma acc atomic capture
+ { x--; v = x; }
+
+}
+
+template<typename T, int I>
+void templ_foo(T v, T x) {
+ // CHECK-NEXT: FunctionTemplateDecl{{.*}}templ_foo
+ // CHECK-NEXT: TemplateTypeParmDecl{{.*}} T
+ // CHECK-NEXT: NonTypeTemplateParmDecl{{.*}} I
+ // CHECK-NEXT: FunctionDecl{{.*}} templ_foo 'void (T, T)'
+ // CHECK-NEXT: ParmVarDecl{{.*}} v 'T'
+ // CHECK-NEXT: ParmVarDecl{{.*}} x 'T'
+ // CHECK-NEXT: CompoundStmt
+
+// CHECK-NEXT: OpenACCAtomicConstruct{{.*}} atomic read
+// CHECK-NEXT: BinaryOperator{{.*}} '<dependent type>' '='
+// CHECK-NEXT: DeclRefExpr{{.*}}'v' 'T'
+// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'T'
+#pragma acc atomic read
+ v = x;
+
+// CHECK-NEXT: OpenACCAtomicConstruct{{.*}} atomic write
+// CHECK-NEXT: BinaryOperator{{.*}} '<dependent type>' '='
+// CHECK-NEXT: DeclRefExpr{{.*}}'v' 'T'
+// CHECK-NEXT: BinaryOperator{{.*}}'<dependent type>' '+'
+// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'T'
+// CHECK-NEXT: DeclRefExpr{{.*}} 'I' 'int'
+#pragma acc atomic write
+ v = x + I;
+
+// CHECK-NEXT: OpenACCAtomicConstruct{{.*}} atomic update
+// CHECK-NEXT: UnaryOperator{{.*}} '<dependent type>' postfix '++'
+// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'T'
+#pragma acc atomic update
+ x++;
+// CHECK-NEXT: OpenACCAtomicConstruct{{.*}} atomic <none>
+// CHECK-NEXT: UnaryOperator{{.*}} '<dependent type>' postfix '--'
+// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'T'
+#pragma acc atomic
+ x--;
+// CHECK-NEXT: OpenACCAtomicConstruct{{.*}} atomic capture
+// CHECK-NEXT: BinaryOperator{{.*}} '<dependent type>' '='
+// CHECK-NEXT: DeclRefExpr{{.*}}'v' 'T'
+// CHECK-NEXT: UnaryOperator{{.*}} '<dependent type>' postfix '++'
+// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'T'
+#pragma acc atomic capture
+ v = x++;
+
+// CHECK-NEXT: OpenACCAtomicConstruct{{.*}} atomic capture
+// CHECK-NEXT: CompoundStmt
+// CHECK-NEXT: UnaryOperator{{.*}} '<dependent type>' postfix '--'
+// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'T'
+// CHECK-NEXT: BinaryOperator{{.*}} '<dependent type>' '='
+// CHECK-NEXT: DeclRefExpr{{.*}}'v' 'T'
+// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'T'
+#pragma acc atomic capture
+ { x--; v = x; }
+
+ // CHECK-NEXT: FunctionDecl{{.*}} templ_foo 'void (int, int)' implicit_instantiation
+ // CHECK-NEXT: TemplateArgument type 'int'
+ // CHECK-NEXT: BuiltinType{{.*}} 'int'
+ // CHECK-NEXT: TemplateArgument integral '5'
+ // CHECK-NEXT: ParmVarDecl{{.*}} v 'int'
+ // CHECK-NEXT: ParmVarDecl{{.*}} x 'int'
+ // CHECK-NEXT: CompoundStmt
+
+// CHECK-NEXT: OpenACCAtomicConstruct{{.*}} atomic read
+// CHECK-NEXT: BinaryOperator{{.*}} 'int' lvalue '='
+// CHECK-NEXT: DeclRefExpr{{.*}}'v' 'int'
+// CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <LValueToRValue>
+// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'int'
+
+// CHECK-NEXT: OpenACCAtomicConstruct{{.*}} atomic write
+// CHECK-NEXT: BinaryOperator{{.*}} 'int' lvalue '='
+// CHECK-NEXT: DeclRefExpr{{.*}}'v' 'int'
+// CHECK-NEXT: BinaryOperator{{.*}}'int' '+'
+// CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <LValueToRValue>
+// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'int'
+// CHECK-NEXT: SubstNonTypeTemplateParmExpr{{.*}} 'int'
+// CHECK-NEXT: NonTypeTemplateParmDecl{{.*}} 'int'{{.*}}I
+// CHECK-NEXT: IntegerLiteral{{.*}} 'int' 5
+
+// CHECK-NEXT: OpenACCAtomicConstruct{{.*}} atomic update
+// CHECK-NEXT: UnaryOperator{{.*}} 'int' postfix '++'
+// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'int'
+
+// CHECK-NEXT: OpenACCAtomicConstruct{{.*}} atomic <none>
+// CHECK-NEXT: UnaryOperator{{.*}} 'int' postfix '--'
+// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'int'
+
+// CHECK-NEXT: OpenACCAtomicConstruct{{.*}} atomic capture
+// CHECK-NEXT: BinaryOperator{{.*}} 'int' lvalue '='
+// CHECK-NEXT: DeclRefExpr{{.*}}'v' 'int'
+// CHECK-NEXT: UnaryOperator{{.*}} 'int' postfix '++'
+// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'int'
+
+// CHECK-NEXT: OpenACCAtomicConstruct{{.*}} atomic capture
+// CHECK-NEXT: CompoundStmt
+// CHECK-NEXT: UnaryOperator{{.*}} 'int' postfix '--'
+// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'int'
+// CHECK-NEXT: BinaryOperator{{.*}} 'int' lvalue '='
+// CHECK-NEXT: DeclRefExpr{{.*}}'v' 'int'
+// CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <LValueToRValue>
+// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'int'
+}
+
+void use() {
+ templ_foo<int, 5>(1, 2);
+}
+#endif
diff --git a/clang/test/SemaOpenACC/atomic-construct.cpp b/clang/test/SemaOpenACC/atomic-construct.cpp
new file mode 100644
index 000000000000000..7357d91d704fd35
--- /dev/null
+++ b/clang/test/SemaOpenACC/atomic-construct.cpp
@@ -0,0 +1,1846 @@
+// RUN: %clang_cc1 %s -fopenacc -Wno-unused-value -verify
+
+void NormalFunc(int I) {
+ // No clauses are valid, but we parse them anyway, just mark them as not valid
+ // on this construct.
+
+ // expected-error at +1{{OpenACC 'copy' clause is not valid on 'atomic' directive}}
+#pragma acc atomic copy(I)
+ I = I + 1;
+ // expected-error at +1{{OpenACC 'copy' clause is not valid on 'atomic' directive}}
+#pragma acc atomic read copy(I)
+ I = I;
+}
+
+struct Struct{
+ Struct *getPtr();
+ Struct &operator++();
+ Struct &operator--();
+ Struct &operator++(int);
+ Struct &operator--(int);
+
+ Struct &operator+=(int);
+ Struct &operator*=(int);
+ Struct &operator-=(int);
+ Struct &operator/=(int);
+ Struct &operator&=(int);
+ Struct &operator|=(int);
+ Struct &operator<<=(int);
+ Struct &operator>>=(int);
+ Struct &operator^=(int);
+ Struct &operator%=(int);
+ Struct &operator!=(int);
+ Struct &operator+();
+ Struct &operator-();
+
+ operator int();
+ void operator()();
+ Struct &operator*();
+ Struct &operator=(int);
+};
+
+int operator+(Struct&, int);
+int operator+(int, Struct&);
+Struct &operator+(Struct&, Struct&);
+Struct &operator*(Struct&, Struct&);
+Struct &operator-(Struct&, Struct&);
+
+Struct S1, S2;
+
+template<typename T>
+T &getRValue();
+
+template<typename T>
+void AtomicReadTemplate(T LHS, T RHS) {
+#pragma acc atomic read
+ LHS = RHS;
+
+ T *LHSPtr, *RHSPtr;
+
+#pragma acc atomic read
+ LHSPtr = RHSPtr;
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic read' directive is invalid}}
+ // expected-note at +2{{right operand to assignment expression must be an l-value}}
+#pragma acc atomic read
+ LHS = RHS + 1;
+
+#pragma acc atomic read
+ *LHSPtr = RHS;
+
+#pragma acc atomic read
+ LHS = *RHSPtr;
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic read' directive is invalid}}
+ // expected-note at +2{{right operand to assignment expression must be an l-value}}
+#pragma acc atomic read
+ LHS = getRValue<T>();
+}
+
+template<typename T>
+void AtomicReadTemplate2(T LHS, T RHS) {
+ // expected-error at +2{{statement associated with OpenACC 'atomic read' directive is invalid}}
+ // expected-note at +2{{left operand to assignment expression must be of scalar type (was 'Struct')}}
+#pragma acc atomic read
+ LHS = RHS;
+
+ T *LHSPtr, *RHSPtr;
+ // Fine, now a pointer.
+#pragma acc atomic read
+ LHSPtr = RHSPtr;
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic read' directive is invalid}}
+ // expected-note at +2{{right operand to assignment expression must be an l-value}}
+#pragma acc atomic read
+ LHS = *RHS.getPtr();
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic read' directive is invalid}}
+ // expected-note at +2{{left operand to assignment expression must be of scalar type (was 'Struct')}}
+#pragma acc atomic read
+ *LHSPtr = RHS;
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic read' directive is invalid}}
+ // expected-note at +2{{left operand to assignment expression must be of scalar type (was 'Struct')}}
+#pragma acc atomic read
+ LHS = *RHSPtr;
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic read' directive is invalid}}
+ // expected-note at +2{{left operand to assignment expression must be an l-value}}
+#pragma acc atomic read
+ getRValue<T>() = getRValue<T>();
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic read' directive is invalid}}
+ // expected-note at +2{{right operand to assignment expression must be an l-value}}
+#pragma acc atomic read
+ LHS = getRValue<T>();
+}
+
+void AtomicRead(int LHS, int RHS) {
+ AtomicReadTemplate(LHS, RHS);
+ AtomicReadTemplate2(S1, S2); // expected-note{{in instantiation of function template specialization}}
+
+#pragma acc atomic read
+ LHS = RHS;
+
+ int *LHSPtr, *RHSPtr;
+
+#pragma acc atomic read
+ LHSPtr = RHSPtr;
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic read' directive is invalid}}
+ // expected-note at +2{{left operand to assignment expression must be of scalar type (was 'Struct')}}
+#pragma acc atomic read
+ S1 = S2;
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic read' directive is invalid}}
+ // expected-note at +2{{right operand to assignment expression must be an l-value}}
+#pragma acc atomic read
+ LHS = RHS + 1;
+
+#pragma acc atomic read
+ *LHSPtr = RHS;
+
+#pragma acc atomic read
+ LHS = *RHSPtr;
+
+ // There is no way to test that = is an overloaded operator, since there
+ // really isn't a way to create an operator= without a class type on one side
+ // or the other.
+}
+
+template<typename T>
+void AtomicWriteTemplate(T LHS, T RHS) {
+#pragma acc atomic write
+ LHS = RHS;
+
+ T *LHSPtr, *RHSPtr;
+#pragma acc atomic write
+ LHSPtr = RHSPtr;
+
+#pragma acc atomic write
+ *LHSPtr = *RHSPtr;
+
+ // allowed, expr is ok.
+#pragma acc atomic write
+ LHS = *RHSPtr;
+
+#pragma acc atomic write
+ LHS = RHS * 2;
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic write' directive is invalid}}
+ // expected-note at +2{{left operand to assignment expression must be an l-value}}
+#pragma acc atomic write
+ getRValue<T>() = getRValue<T>();
+
+#pragma acc atomic write
+ LHS = getRValue<T>();
+}
+
+template<typename T>
+void AtomicWriteTemplate2(T LHS, T RHS) {
+ // expected-error at +2{{statement associated with OpenACC 'atomic write' directive is invalid}}
+ // expected-note at +2{{left operand to assignment expression must be of scalar type (was 'Struct')}}
+#pragma acc atomic write
+ LHS = RHS;
+
+ T *LHSPtr, *RHSPtr;
+#pragma acc atomic write
+ LHSPtr = RHSPtr;
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic write' directive is invalid}}
+ // expected-note at +2{{left operand to assignment expression must be of scalar type (was 'Struct')}}
+#pragma acc atomic write
+ LHS = *RHSPtr;
+
+#pragma acc atomic write
+ LHSPtr = RHS.getPtr();
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic write' directive is invalid}}
+ // expected-note at +2{{left operand to assignment expression must be an l-value}}
+#pragma acc atomic write
+ getRValue<T>() = getRValue<T>();
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic write' directive is invalid}}
+ // expected-note at +2{{left operand to assignment expression must be of scalar type (was 'Struct')}}
+#pragma acc atomic write
+ LHS = getRValue<T>();
+}
+
+void AtomicWrite(int LHS, int RHS) {
+ AtomicWriteTemplate(LHS, RHS);
+ AtomicWriteTemplate2(S1, S2); // expected-note{{in instantiation of function template specialization}}
+
+#pragma acc atomic write
+ LHS = RHS;
+
+ int *LHSPtr, *RHSPtr;
+#pragma acc atomic write
+ LHSPtr = RHSPtr;
+
+#pragma acc atomic write
+ *LHSPtr = *RHSPtr;
+
+ // allowed, expr is ok.
+#pragma acc atomic write
+ LHS = *RHSPtr;
+
+#pragma acc atomic write
+ LHS = RHS * 2;
+}
+
+template<typename T>
+void AtomicUpdateTemplate(T LHS, T RHS) {
+#pragma acc atomic
+ LHS++;
+
+#pragma acc atomic update
+ LHS--;
+
+#pragma acc atomic
+ ++LHS;
+
+#pragma acc atomic update
+ --LHS;
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic' directive is invalid}}
+ // expected-note at +2{{unary operator not supported, only increment and decrement operations permitted}}
+#pragma acc atomic
+ +LHS;
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic update' directive is invalid}}
+ // expected-note at +2{{unary operator not supported, only increment and decrement operations permitted}}
+#pragma acc atomic update
+ -LHS;
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic update' directive is invalid}}
+ // expected-note at +2{{expected binary operation on right hand side of assignment operator}}
+#pragma acc atomic update
+ LHS = RHS;
+
+ T *LHSPtr, *RHSPtr;
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic' directive is invalid}}
+ // expected-note at +2{{expected binary operation on right hand side of assignment operator}}
+#pragma acc atomic
+ *LHSPtr = *RHSPtr;
+
+ // x binop= expr;
+#pragma acc atomic
+ LHS += 1 + RHS;
+#pragma acc atomic update
+ LHS *= 1 + RHS;
+#pragma acc atomic
+ LHS -= 1 + RHS;
+#pragma acc atomic update
+ LHS /= 1 + RHS;
+#pragma acc atomic
+ LHS &= 1 + RHS;
+#pragma acc atomic update
+ LHS ^= 1 + RHS;
+#pragma acc atomic
+ LHS |= 1 + RHS;
+#pragma acc atomic update
+ LHS <<= 1 + RHS;
+#pragma acc atomic
+ LHS >>= 1 + RHS;
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic update' directive is invalid}}
+ // expected-note at +2{{compound binary operator not supported, only +=, *=, -=, /=, &=, ^=, |=, <<=, or >>= are permitted}}
+#pragma acc atomic update
+ LHS != 1 + RHS;
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic update' directive is invalid}}
+ // expected-note at +2{{compound binary operator not supported, only +=, *=, -=, /=, &=, ^=, |=, <<=, or >>= are permitted}}
+#pragma acc atomic update
+ LHS <= 1 + RHS;
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic update' directive is invalid}}
+ // expected-note at +2{{compound binary operator not supported, only +=, *=, -=, /=, &=, ^=, |=, <<=, or >>= are permitted}}
+#pragma acc atomic update
+ LHS >= 1 + RHS;
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic update' directive is invalid}}
+ // expected-note at +2{{compound binary operator not supported, only +=, *=, -=, /=, &=, ^=, |=, <<=, or >>= are permitted}}
+#pragma acc atomic update
+ LHS %= 1 + RHS;
+
+ // x = x binop expr.
+#pragma acc atomic
+ LHS = LHS + getRValue<T>();
+#pragma acc atomic update
+ LHS = LHS * getRValue<T>();
+#pragma acc atomic update
+ LHS = LHS - getRValue<T>();
+#pragma acc atomic update
+ LHS = LHS / getRValue<T>();
+#pragma acc atomic update
+ LHS = LHS & getRValue<T>();
+#pragma acc atomic update
+ LHS = LHS ^ getRValue<T>();
+#pragma acc atomic update
+ LHS = LHS | getRValue<T>();
+#pragma acc atomic update
+ LHS = LHS << getRValue<T>();
+#pragma acc atomic update
+ LHS = LHS >> getRValue<T>();
+ // expected-error at +2{{statement associated with OpenACC 'atomic update' directive is invalid}}
+ // expected-note at +2{{binary operator not supported, only +, *, -, /, &, ^, |, <<, or >> are permitted}}
+#pragma acc atomic update
+ LHS = LHS < getRValue<T>();
+ // expected-error at +2{{statement associated with OpenACC 'atomic update' directive is invalid}}
+ // expected-note at +2{{binary operator not supported, only +, *, -, /, &, ^, |, <<, or >> are permitted}}
+#pragma acc atomic update
+ LHS = LHS > getRValue<T>();
+ // expected-error at +2{{statement associated with OpenACC 'atomic update' directive is invalid}}
+ // expected-note at +2{{binary operator not supported, only +, *, -, /, &, ^, |, <<, or >> are permitted}}
+#pragma acc atomic update
+ LHS = LHS <= getRValue<T>();
+ // expected-error at +2{{statement associated with OpenACC 'atomic update' directive is invalid}}
+ // expected-note at +2{{binary operator not supported, only +, *, -, /, &, ^, |, <<, or >> are permitted}}
+#pragma acc atomic update
+ LHS = LHS >= getRValue<T>();
+#pragma acc atomic update
+ LHS = LHS ^ getRValue<T>();
+
+
+ // x = expr binop x.
+#pragma acc atomic
+ LHS = getRValue<T>() + LHS;
+#pragma acc atomic update
+ LHS = getRValue<T>() * LHS;
+#pragma acc atomic update
+ LHS = getRValue<T>() - LHS;
+#pragma acc atomic update
+ LHS = getRValue<T>() / LHS;
+#pragma acc atomic update
+ LHS = getRValue<T>() & LHS;
+#pragma acc atomic update
+ LHS = getRValue<T>() ^ LHS;
+#pragma acc atomic update
+ LHS = getRValue<T>() | LHS;
+#pragma acc atomic update
+ LHS = getRValue<T>() << LHS;
+#pragma acc atomic update
+ LHS = getRValue<T>() >> LHS;
+ // expected-error at +2{{statement associated with OpenACC 'atomic update' directive is invalid}}
+ // expected-note at +2{{binary operator not supported, only +, *, -, /, &, ^, |, <<, or >> are permitted}}
+#pragma acc atomic update
+ LHS = getRValue<T>() < LHS;
+ // expected-error at +2{{statement associated with OpenACC 'atomic update' directive is invalid}}
+ // expected-note at +2{{binary operator not supported, only +, *, -, /, &, ^, |, <<, or >> are permitted}}
+#pragma acc atomic update
+ LHS = getRValue<T>() > LHS;
+ // expected-error at +2{{statement associated with OpenACC 'atomic update' directive is invalid}}
+ // expected-note at +2{{binary operator not supported, only +, *, -, /, &, ^, |, <<, or >> are permitted}}
+#pragma acc atomic update
+ LHS = getRValue<T>() <= LHS;
+ // expected-error at +2{{statement associated with OpenACC 'atomic update' directive is invalid}}
+ // expected-note at +2{{binary operator not supported, only +, *, -, /, &, ^, |, <<, or >> are permitted}}
+#pragma acc atomic update
+ LHS = getRValue<T>() >= LHS;
+#pragma acc atomic update
+ LHS = getRValue<T>() ^ LHS;
+
+#pragma acc atomic update
+ LHS = LHS + getRValue<T>();
+ // expected-error at +2{{statement associated with OpenACC 'atomic update' directive is invalid}}
+ // expected-note at +2{{left hand side of assignment operation('LHS') must match one side of the sub-operation on the right hand side('RHS' and 'getRValue<T>()')}}
+#pragma acc atomic update
+ LHS = RHS + getRValue<T>();
+
+#pragma acc atomic update
+ LHS = getRValue<T>() - LHS;
+ // expected-error at +2{{statement associated with OpenACC 'atomic update' directive is invalid}}
+ // expected-note at +2{{left hand side of assignment operation('LHS') must match one side of the sub-operation on the right hand side('getRValue<T>()' and 'RHS')}}
+#pragma acc atomic update
+ LHS = getRValue<T>() + RHS;
+}
+
+template<typename T>
+void AtomicUpdateTemplate2(T LHS, T RHS) {
+ // expected-error at +2{{statement associated with OpenACC 'atomic' directive is invalid}}
+ // expected-note at +2{{operand to increment expression must be of scalar type (was 'Struct')}}
+#pragma acc atomic
+ LHS++;
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic update' directive is invalid}}
+ // expected-note at +2{{operand to decrement expression must be of scalar type (was 'Struct')}}
+#pragma acc atomic update
+ LHS--;
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic' directive is invalid}}
+ // expected-note at +2{{operand to increment expression must be of scalar type (was 'Struct')}}
+#pragma acc atomic
+ ++LHS;
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic update' directive is invalid}}
+ // expected-note at +2{{operand to decrement expression must be of scalar type (was 'Struct')}}
+#pragma acc atomic update
+ --LHS;
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic' directive is invalid}}
+ // expected-note at +2{{unary operator not supported, only increment and decrement operations permitted}}
+#pragma acc atomic
+ +LHS;
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic update' directive is invalid}}
+ // expected-note at +2{{unary operator not supported, only increment and decrement operations permitted}}
+#pragma acc atomic update
+ -LHS;
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic' directive is invalid}}
+ // expected-note at +2{{expected assignment, compound assignment, increment, or decrement expression}}
+#pragma acc atomic
+ LHS();
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic' directive is invalid}}
+ // expected-note at +2{{unary operator not supported, only increment and decrement operations permitted}}
+#pragma acc atomic
+ *LHS;
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic update' directive is invalid}}
+ // expected-note at +2{{expected binary operation on right hand side of assignment operator}}
+#pragma acc atomic update
+ LHS = RHS;
+
+ T *LHSPtr, *RHSPtr;
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic' directive is invalid}}
+ // expected-note at +2{{expected binary operation on right hand side of assignment operator}}
+#pragma acc atomic
+ *LHSPtr = *RHSPtr;
+
+ // x binop= expr;
+ // expected-error at +2{{statement associated with OpenACC 'atomic' directive is invalid}}
+ // expected-note at +2{{left operand to compound assignment expression must be of scalar type (was 'Struct')}}
+#pragma acc atomic
+ LHS += 1 + RHS;
+ // expected-error at +2{{statement associated with OpenACC 'atomic update' directive is invalid}}
+ // expected-note at +2{{left operand to compound assignment expression must be of scalar type (was 'Struct')}}
+#pragma acc atomic update
+ LHS *= 1 + RHS;
+ // expected-error at +2{{statement associated with OpenACC 'atomic' directive is invalid}}
+ // expected-note at +2{{left operand to compound assignment expression must be of scalar type (was 'Struct')}}
+#pragma acc atomic
+ LHS -= 1 + RHS;
+ // expected-error at +2{{statement associated with OpenACC 'atomic update' directive is invalid}}
+ // expected-note at +2{{left operand to compound assignment expression must be of scalar type (was 'Struct')}}
+#pragma acc atomic update
+ LHS /= 1 + RHS;
+ // expected-error at +2{{statement associated with OpenACC 'atomic' directive is invalid}}
+ // expected-note at +2{{left operand to compound assignment expression must be of scalar type (was 'Struct')}}
+#pragma acc atomic
+ LHS &= 1 + RHS;
+ // expected-error at +2{{statement associated with OpenACC 'atomic update' directive is invalid}}
+ // expected-note at +2{{left operand to compound assignment expression must be of scalar type (was 'Struct')}}
+#pragma acc atomic update
+ LHS ^= 1 + RHS;
+ // expected-error at +2{{statement associated with OpenACC 'atomic' directive is invalid}}
+ // expected-note at +2{{left operand to compound assignment expression must be of scalar type (was 'Struct')}}
+#pragma acc atomic
+ LHS |= 1 + RHS;
+ // expected-error at +2{{statement associated with OpenACC 'atomic update' directive is invalid}}
+ // expected-note at +2{{left operand to compound assignment expression must be of scalar type (was 'Struct')}}
+#pragma acc atomic update
+ LHS <<= 1 + RHS;
+ // expected-error at +2{{statement associated with OpenACC 'atomic' directive is invalid}}
+ // expected-note at +2{{left operand to compound assignment expression must be of scalar type (was 'Struct')}}
+#pragma acc atomic
+ LHS >>= 1 + RHS;
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic update' directive is invalid}}
+ // expected-note at +2{{compound binary operator not supported, only +=, *=, -=, /=, &=, ^=, |=, <<=, or >>= are permitted}}
+#pragma acc atomic update
+ LHS != 1 + RHS;
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic update' directive is invalid}}
+ // expected-note at +2{{compound binary operator not supported, only +=, *=, -=, /=, &=, ^=, |=, <<=, or >>= are permitted}}
+#pragma acc atomic update
+ LHS %= 1 + RHS;
+
+ // x = x binop expr.
+ // expected-error at +2{{statement associated with OpenACC 'atomic' directive is invalid}}
+ // expected-note at +2{{left operand to assignment expression must be of scalar type (was 'Struct')}}
+#pragma acc atomic
+ LHS = LHS + getRValue<T>();
+ // expected-error at +2{{statement associated with OpenACC 'atomic update' directive is invalid}}
+ // expected-note at +2{{left operand to assignment expression must be of scalar type (was 'Struct')}}
+#pragma acc atomic update
+ LHS = LHS * getRValue<T>();
+ // expected-error at +2{{statement associated with OpenACC 'atomic update' directive is invalid}}
+ // expected-note at +2{{left operand to assignment expression must be of scalar type (was 'Struct')}}
+#pragma acc atomic update
+ LHS = LHS - getRValue<T>();
+
+ // x = expr binop x.
+ // expected-error at +2{{statement associated with OpenACC 'atomic' directive is invalid}}
+ // expected-note at +2{{left operand to assignment expression must be of scalar type (was 'Struct')}}
+#pragma acc atomic
+ LHS = getRValue<T>() + LHS;
+ // expected-error at +2{{statement associated with OpenACC 'atomic update' directive is invalid}}
+ // expected-note at +2{{left operand to assignment expression must be of scalar type (was 'Struct')}}
+#pragma acc atomic update
+ LHS = getRValue<T>() * LHS;
+ // expected-error at +2{{statement associated with OpenACC 'atomic update' directive is invalid}}
+ // expected-note at +2{{left operand to assignment expression must be of scalar type (was 'Struct')}}
+#pragma acc atomic update
+ LHS = getRValue<T>() - LHS;
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic update' directive is invalid}}
+ // expected-note at +2{{left operand to assignment expression must be of scalar type (was 'Struct')}}
+#pragma acc atomic update
+ LHS = LHS + getRValue<T>();
+ // expected-error at +2{{statement associated with OpenACC 'atomic update' directive is invalid}}
+ // expected-note at +2{{left hand side of assignment operation('LHS') must match one side of the sub-operation on the right hand side('RHS' and 'getRValue<T>()')}}
+#pragma acc atomic update
+ LHS = RHS + getRValue<T>();
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic update' directive is invalid}}
+ // expected-note at +2{{left operand to assignment expression must be of scalar type (was 'Struct')}}
+#pragma acc atomic update
+ LHS = getRValue<T>() - LHS;
+ // expected-error at +2{{statement associated with OpenACC 'atomic update' directive is invalid}}
+ // expected-note at +2{{left hand side of assignment operation('LHS') must match one side of the sub-operation on the right hand side('getRValue<T>()' and 'RHS')}}
+#pragma acc atomic update
+ LHS = getRValue<T>() + RHS;
+}
+
+void AtomicUpdate() {
+ AtomicUpdateTemplate(1, 2);
+ AtomicUpdateTemplate2(S1, S2); //expected-note{{in instantiation of function template specialization}}
+
+ int I, J;
+
+#pragma acc atomic
+ I++;
+#pragma acc atomic update
+ --I;
+ // expected-error at +2{{statement associated with OpenACC 'atomic' directive is invalid}}
+ // expected-note at +2{{operand to increment expression must be of scalar type (was 'Struct')}}
+#pragma acc atomic
+ S1++;
+ // expected-error at +2{{statement associated with OpenACC 'atomic update' directive is invalid}}
+ // expected-note at +2{{operand to decrement expression must be of scalar type (was 'Struct')}}
+#pragma acc atomic update
+ --S2;
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic' directive is invalid}}
+ // expected-note at +2{{unary operator not supported, only increment and decrement operations permitted}}
+#pragma acc atomic
+ +I;
+ // expected-error at +2{{statement associated with OpenACC 'atomic update' directive is invalid}}
+ // expected-note at +2{{unary operator not supported, only increment and decrement operations permitted}}
+#pragma acc atomic update
+ -J;
+
+#pragma acc atomic update
+ I ^= 1 + J;
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic update' directive is invalid}}
+ // expected-note at +2{{compound binary operator not supported, only +=, *=, -=, /=, &=, ^=, |=, <<=, or >>= are permitted}}
+#pragma acc atomic update
+ I%= 1 + J;
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic update' directive is invalid}}
+ // expected-note at +2{{left operand to compound assignment expression must be of scalar type (was 'Struct')}}
+#pragma acc atomic update
+ S1 ^= 1 + J;
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic update' directive is invalid}}
+ // expected-note at +2{{compound binary operator not supported, only +=, *=, -=, /=, &=, ^=, |=, <<=, or >>= are permitted}}
+#pragma acc atomic update
+ S2 %= 1 + J;
+
+#pragma acc atomic update
+ I = I + getRValue<int>();
+ // expected-error at +2{{statement associated with OpenACC 'atomic update' directive is invalid}}
+ // expected-note at +2{{left hand side of assignment operation('I') must match one side of the sub-operation on the right hand side('J' and 'getRValue<int>()')}}
+#pragma acc atomic update
+ I = J + getRValue<int>();
+
+#pragma acc atomic update
+ I = getRValue<int>() - I;
+ // expected-error at +2{{statement associated with OpenACC 'atomic update' directive is invalid}}
+ // expected-note at +2{{left hand side of assignment operation('I') must match one side of the sub-operation on the right hand side('getRValue<int>()' and 'J')}}
+#pragma acc atomic update
+ I = getRValue<int>() + J;
+}
+
+template<typename T>
+void AtomicCaptureTemplateSimple(T LHS, T RHS) {
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{expected assignment expression}}
+#pragma acc atomic capture
+ LHS++;
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{expected assignment expression}}
+#pragma acc atomic capture
+--LHS;
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{expected assignment expression}}
+#pragma acc atomic capture
+ LHS += 1 + RHS;
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{expected assignment, compound assignment, increment, or decrement expression}}
+#pragma acc atomic capture
+ LHS = RHS;
+
+#pragma acc atomic capture
+ LHS = RHS++;
+
+#pragma acc atomic capture
+ LHS = RHS--;
+
+#pragma acc atomic capture
+ LHS = ++RHS;
+
+#pragma acc atomic capture
+ LHS = --RHS;
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{unary operator not supported, only increment and decrement operations permitted}}
+#pragma acc atomic capture
+ LHS = +RHS;
+
+#pragma acc atomic capture
+ LHS = RHS += 1 + RHS;
+#pragma acc atomic capture
+ LHS = RHS *= 1 + RHS;
+#pragma acc atomic capture
+ LHS = RHS -= 1 + RHS;
+#pragma acc atomic capture
+ LHS = RHS /= 1 + RHS;
+#pragma acc atomic capture
+ LHS = RHS &= 1 + RHS;
+#pragma acc atomic capture
+ LHS = RHS ^= 1 + RHS;
+#pragma acc atomic capture
+ LHS = RHS >>= 1 + RHS;
+#pragma acc atomic capture
+ LHS = RHS |= 1 + RHS;
+#pragma acc atomic capture
+ LHS = RHS <<= 1 + RHS;
+#pragma acc atomic capture
+ LHS = RHS >>= 1 + RHS;
+
+#pragma acc atomic capture
+ LHS = RHS ^= 1 + RHS;
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{compound binary operator not supported, only +=, *=, -=, /=, &=, ^=, |=, <<=, or >>= are permitted}}
+#pragma acc atomic capture
+ LHS = RHS <= 1 + RHS;
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{compound binary operator not supported, only +=, *=, -=, /=, &=, ^=, |=, <<=, or >>= are permitted}}
+#pragma acc atomic capture
+ LHS = RHS >= 1 + RHS;
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{compound binary operator not supported, only +=, *=, -=, /=, &=, ^=, |=, <<=, or >>= are permitted}}
+#pragma acc atomic capture
+ LHS = RHS + 1;
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{compound binary operator not supported, only +=, *=, -=, /=, &=, ^=, |=, <<=, or >>= are permitted}}
+#pragma acc atomic capture
+ LHS = RHS < 1 + RHS;
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{compound binary operator not supported, only +=, *=, -=, /=, &=, ^=, |=, <<=, or >>= are permitted}}
+#pragma acc atomic capture
+ LHS = RHS > 1 + RHS;
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{compound binary operator not supported, only +=, *=, -=, /=, &=, ^=, |=, <<=, or >>= are permitted}}
+#pragma acc atomic capture
+ LHS = RHS ^ 1 + RHS;
+
+#pragma acc atomic capture
+ LHS = RHS = RHS + 1;
+#pragma acc atomic capture
+ LHS = RHS = 1 + RHS;
+#pragma acc atomic capture
+ LHS = RHS = RHS * 1;
+#pragma acc atomic capture
+ LHS = RHS = 1 * RHS;
+#pragma acc atomic capture
+ LHS = RHS = RHS / 1;
+#pragma acc atomic capture
+ LHS = RHS = 1 / RHS;
+#pragma acc atomic capture
+ LHS = RHS = RHS ^ 1;
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{binary operator not supported, only +, *, -, /, &, ^, |, <<, or >> are permitted}}
+#pragma acc atomic capture
+ LHS = RHS = 1 % RHS;
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{binary operator not supported, only +, *, -, /, &, ^, |, <<, or >> are permitted}}
+#pragma acc atomic capture
+ LHS = RHS = RHS < 1;
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{binary operator not supported, only +, *, -, /, &, ^, |, <<, or >> are permitted}}
+#pragma acc atomic capture
+ LHS = RHS = 1 > RHS;
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{left hand side of assignment operation('LHS') must match one side of the sub-operation on the right hand side('RHS' and 'getRValue<T>()')}}
+#pragma acc atomic capture
+ LHS = LHS = RHS + getRValue<T>();
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{left hand side of assignment operation('LHS') must match one side of the sub-operation on the right hand side('getRValue<T>()' and 'RHS')}}
+#pragma acc atomic capture
+ LHS = LHS = getRValue<T>() + RHS;
+}
+template<typename T>
+void AtomicCaptureTemplateSimple2(T LHS, T RHS) {
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{expected assignment expression}}
+#pragma acc atomic capture
+ LHS++;
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{expected assignment expression}}
+#pragma acc atomic capture
+--LHS;
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{expected assignment expression}}
+#pragma acc atomic capture
+ LHS += 1 + RHS;
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{expected assignment, compound assignment, increment, or decrement expression}}
+#pragma acc atomic capture
+ LHS = RHS;
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{left operand to assignment expression must be of scalar type (was 'Struct')}}
+#pragma acc atomic capture
+ LHS = RHS++;
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{left operand to assignment expression must be of scalar type (was 'Struct')}}
+#pragma acc atomic capture
+ LHS = RHS--;
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{left operand to assignment expression must be of scalar type (was 'Struct')}}
+#pragma acc atomic capture
+ LHS = ++RHS;
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{left operand to assignment expression must be of scalar type (was 'Struct')}}
+#pragma acc atomic capture
+ LHS = --RHS;
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{unary operator not supported, only increment and decrement operations permitted}}
+#pragma acc atomic capture
+ LHS = +RHS;
+
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{left operand to assignment expression must be of scalar type (was 'Struct')}}
+#pragma acc atomic capture
+ LHS = RHS += 1 + RHS;
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{left operand to assignment expression must be of scalar type (was 'Struct')}}
+#pragma acc atomic capture
+ LHS = RHS *= 1 + RHS;
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{left operand to assignment expression must be of scalar type (was 'Struct')}}
+#pragma acc atomic capture
+ LHS = RHS -= 1 + RHS;
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{left operand to assignment expression must be of scalar type (was 'Struct')}}
+#pragma acc atomic capture
+ LHS = RHS /= 1 + RHS;
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{left operand to assignment expression must be of scalar type (was 'Struct')}}
+#pragma acc atomic capture
+ LHS = RHS &= 1 + RHS;
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{left operand to assignment expression must be of scalar type (was 'Struct')}}
+#pragma acc atomic capture
+ LHS = RHS ^= 1 + RHS;
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{left operand to assignment expression must be of scalar type (was 'Struct')}}
+#pragma acc atomic capture
+ LHS = RHS >>= 1 + RHS;
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{left operand to assignment expression must be of scalar type (was 'Struct')}}
+#pragma acc atomic capture
+ LHS = RHS |= 1 + RHS;
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{left operand to assignment expression must be of scalar type (was 'Struct')}}
+#pragma acc atomic capture
+ LHS = RHS <<= 1 + RHS;
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{left operand to assignment expression must be of scalar type (was 'Struct')}}
+#pragma acc atomic capture
+ LHS = RHS >>= 1 + RHS;
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{left operand to assignment expression must be of scalar type (was 'Struct')}}
+#pragma acc atomic capture
+ LHS = RHS ^= 1 + RHS;
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{compound binary operator not supported, only +=, *=, -=, /=, &=, ^=, |=, <<=, or >>= are permitted}}
+#pragma acc atomic capture
+ LHS = RHS <= 1 + RHS;
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{compound binary operator not supported, only +=, *=, -=, /=, &=, ^=, |=, <<=, or >>= are permitted}}
+#pragma acc atomic capture
+ LHS = RHS >= 1 + RHS;
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{compound binary operator not supported, only +=, *=, -=, /=, &=, ^=, |=, <<=, or >>= are permitted}}
+#pragma acc atomic capture
+ LHS = RHS + 1;
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{compound binary operator not supported, only +=, *=, -=, /=, &=, ^=, |=, <<=, or >>= are permitted}}
+#pragma acc atomic capture
+ LHS = RHS < 1 + RHS;
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{compound binary operator not supported, only +=, *=, -=, /=, &=, ^=, |=, <<=, or >>= are permitted}}
+#pragma acc atomic capture
+ LHS = RHS > 1 + RHS;
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{compound binary operator not supported, only +=, *=, -=, /=, &=, ^=, |=, <<=, or >>= are permitted}}
+#pragma acc atomic capture
+ LHS = RHS ^ 1 + RHS;
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{left operand to assignment expression must be of scalar type (was 'Struct')}}
+#pragma acc atomic capture
+ LHS = RHS = RHS + 1;
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{left operand to assignment expression must be of scalar type (was 'Struct')}}
+#pragma acc atomic capture
+ LHS = RHS = 1 + RHS;
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{left operand to assignment expression must be of scalar type (was 'Struct')}}
+#pragma acc atomic capture
+ LHS = RHS = RHS * 1;
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{left operand to assignment expression must be of scalar type (was 'Struct')}}
+#pragma acc atomic capture
+ LHS = RHS = 1 * RHS;
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{left operand to assignment expression must be of scalar type (was 'Struct')}}
+#pragma acc atomic capture
+ LHS = RHS = RHS / 1;
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{left operand to assignment expression must be of scalar type (was 'Struct')}}
+#pragma acc atomic capture
+ LHS = RHS = 1 / RHS;
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{left operand to assignment expression must be of scalar type (was 'Struct')}}
+#pragma acc atomic capture
+ LHS = RHS = RHS ^ 1;
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{binary operator not supported, only +, *, -, /, &, ^, |, <<, or >> are permitted}}
+#pragma acc atomic capture
+ LHS = RHS = 1 % RHS;
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{binary operator not supported, only +, *, -, /, &, ^, |, <<, or >> are permitted}}
+#pragma acc atomic capture
+ LHS = RHS = RHS < 1;
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{binary operator not supported, only +, *, -, /, &, ^, |, <<, or >> are permitted}}
+#pragma acc atomic capture
+ LHS = RHS = 1 > RHS;
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{left hand side of assignment operation('LHS') must match one side of the sub-operation on the right hand side('RHS' and 'getRValue<T>()')}}
+#pragma acc atomic capture
+ LHS = LHS = RHS + getRValue<T>();
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{left hand side of assignment operation('LHS') must match one side of the sub-operation on the right hand side('getRValue<T>()' and 'RHS')}}
+#pragma acc atomic capture
+ LHS = LHS = getRValue<T>() + RHS;
+}
+
+void AtomicCaptureSimple(int LHS, int RHS) {
+ AtomicCaptureTemplateSimple(1, 2);
+ AtomicCaptureTemplateSimple2(S1, S2); //expected-note{{in instantiation of function template specialization}}
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{expected assignment expression}}
+#pragma acc atomic capture
+ LHS++;
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{expected assignment expression}}
+#pragma acc atomic capture
+--LHS;
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{expected assignment expression}}
+#pragma acc atomic capture
+ LHS += 1 + RHS;
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{expected assignment, compound assignment, increment, or decrement expression}}
+#pragma acc atomic capture
+ LHS = RHS;
+
+#pragma acc atomic capture
+ LHS = RHS++;
+
+#pragma acc atomic capture
+ LHS = RHS--;
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{left operand to assignment expression must be of scalar type (was 'Struct')}}
+#pragma acc atomic capture
+ S1 = ++S2;
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{left operand to assignment expression must be of scalar type (was 'Struct')}}
+#pragma acc atomic capture
+ S1 = --S2 ;
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{unary operator not supported, only increment and decrement operations permitted}}
+#pragma acc atomic capture
+ LHS = +RHS;
+
+#pragma acc atomic capture
+ LHS = RHS += 1 + RHS;
+#pragma acc atomic capture
+ LHS = RHS *= 1 + RHS;
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{left operand to assignment expression must be of scalar type (was 'Struct')}}
+#pragma acc atomic capture
+ S1 = RHS -= 1 + RHS;
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{expected assignment, compound assignment, increment, or decrement expression}}
+#pragma acc atomic capture
+ LHS = S1 /= 1 + RHS;
+#pragma acc atomic capture
+ LHS = RHS &= 1 + S2;
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{expected assignment, compound assignment, increment, or decrement expression}}
+#pragma acc atomic capture
+ LHS = S1^= 1 + S2;
+
+#pragma acc atomic capture
+ LHS = RHS ^= 1 + RHS;
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{compound binary operator not supported, only +=, *=, -=, /=, &=, ^=, |=, <<=, or >>= are permitted}}
+#pragma acc atomic capture
+ LHS = RHS <= 1 + RHS;
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{left operand to assignment expression must be of scalar type (was 'Struct')}}
+#pragma acc atomic capture
+ S1 = RHS ^= 1 + RHS;
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{compound binary operator not supported, only +=, *=, -=, /=, &=, ^=, |=, <<=, or >>= are permitted}}
+#pragma acc atomic capture
+ LHS = S1 <= 1 + RHS;
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{compound binary operator not supported, only +=, *=, -=, /=, &=, ^=, |=, <<=, or >>= are permitted}}
+#pragma acc atomic capture
+ LHS = RHS <= 1 + S2;
+
+#pragma acc atomic capture
+ LHS = RHS = RHS + 1;
+#pragma acc atomic capture
+ LHS = RHS = 1 + RHS;
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{left operand to assignment expression must be of scalar type (was 'Struct')}}
+#pragma acc atomic capture
+ S1 = RHS = RHS * 1;
+ // A little weird, because this contains a 'operator int' call here rather
+ // than a conversion, so the diagnostic could be better.
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{expected assignment, compound assignment, increment, or decrement expression}}
+#pragma acc atomic capture
+ LHS = S2 = 1 * S2;
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{binary operator not supported, only +, *, -, /, &, ^, |, <<, or >> are permitted}}
+#pragma acc atomic capture
+ LHS = RHS = RHS < 1;
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{binary operator not supported, only +, *, -, /, &, ^, |, <<, or >> are permitted}}
+#pragma acc atomic capture
+ LHS = RHS = 1 > RHS;
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{left operand to assignment expression must be of scalar type (was 'Struct')}}
+#pragma acc atomic capture
+ S1 = RHS = RHS < 1;
+
+ // A little weird, because this contains a 'operator int' call here rather
+ // than a conversion, so the diagnostic could be better.
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{expected assignment, compound assignment, increment, or decrement expression}}
+#pragma acc atomic capture
+ LHS = S1 = 1 > S1;
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{left hand side of assignment operation('LHS') must match one side of the sub-operation on the right hand side('RHS' and 'getRValue<int>()')}}
+#pragma acc atomic capture
+ LHS = LHS = RHS + getRValue<int>();
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{left hand side of assignment operation('LHS') must match one side of the sub-operation on the right hand side('getRValue<int>()' and 'RHS')}}
+#pragma acc atomic capture
+ LHS = LHS = getRValue<int>() + RHS;
+}
+
+template<typename T>
+void AtomicCaptureTemplateCompound(T LHS, T RHS) {
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{expected assignment, compound assignment, increment, or decrement expression}}
+#pragma acc atomic capture
+ {
+ }
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +4{{expected assignment, compound assignment, increment, or decrement expression}}
+#pragma acc atomic capture
+ {
+ LHS = RHS;
+ }
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +3{{'atomic capture' with a compound statement only supports two statements}}
+#pragma acc atomic capture
+ {
+ LHS = RHS; RHS += 1; LHS=RHS;
+ }
+
+
+#pragma acc atomic capture
+ {
+ LHS++;
+ RHS = LHS;
+ }
+
+#pragma acc atomic capture
+ {
+ ++LHS;
+ RHS = LHS;
+ }
+
+#pragma acc atomic capture
+ {
+ --LHS;
+ RHS = LHS;
+ }
+
+
+#pragma acc atomic capture
+ {
+ LHS--;
+ // expected-error at -3{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{variable on right hand side of assignment('RHS') must match variable used in unary expression('LHS') from the first statement}}
+ LHS = RHS;
+ }
+
+#pragma acc atomic capture
+ {
+ // expected-error at -2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{unary operator not supported, only increment and decrement operations permitted}}
+ -LHS;
+ RHS = LHS;
+ }
+
+#pragma acc atomic capture
+ {
+ --LHS;
+ // expected-error at -3{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{expected assignment expression}}
+ RHS += LHS;
+ }
+
+ // { x binop = expr; v = x; }
+#pragma acc atomic capture
+ {
+ LHS += 1;
+ RHS = LHS;
+ }
+#pragma acc atomic capture
+ {
+ LHS *= 1;
+ // expected-error at -3{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{variable on right hand side of assignment('RHS') must match variable used on left hand side of compound assignment('LHS') from the first statement}}
+ LHS = RHS;
+ }
+#pragma acc atomic capture
+ {
+ LHS /= 1;
+ // expected-error at -3{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{expected assignment expression}}
+ RHS += LHS;
+ }
+
+ // { x = x binop expr; v = x; }
+#pragma acc atomic capture
+ {
+ LHS = LHS + 1;
+ RHS = LHS;
+ }
+
+#pragma acc atomic capture
+ {
+ // expected-error at -2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{left hand side of assignment operation('LHS') must match one side of the sub-operation on the right hand side('RHS' and '1')}}
+ LHS = RHS - 1;
+ RHS = LHS;
+ }
+#pragma acc atomic capture
+ {
+ LHS = LHS * 1;
+ // expected-error at -3{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{variable on right hand side of assignment('RHS') must match variable used on left hand side of assignment('LHS') from the first statement}}
+ RHS = RHS;
+ }
+#pragma acc atomic capture
+ {
+ LHS = LHS / 1;
+ // expected-error at -3{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{expected assignment expression}}
+ RHS += LHS;
+ }
+
+ // { x = expr binop x; v = x; }
+#pragma acc atomic capture
+ {
+ LHS = 1 ^ LHS;
+ RHS = LHS;
+ }
+
+#pragma acc atomic capture
+ {
+ // expected-error at -2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{left hand side of assignment operation('LHS') must match one side of the sub-operation on the right hand side('1' and 'RHS')}}
+ LHS = 1 & RHS;
+ RHS = LHS;
+ }
+#pragma acc atomic capture
+ {
+ LHS = LHS | 1;
+ // expected-error at -3{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{variable on right hand side of assignment('RHS') must match variable used on left hand side of assignment('LHS') from the first statement}}
+ RHS = RHS;
+ }
+#pragma acc atomic capture
+ {
+ LHS = LHS << 1;
+ // expected-error at -3{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{expected assignment expression}}
+ RHS += LHS;
+ }
+
+ // { v = x; x binop = expr; }
+#pragma acc atomic capture
+ {
+ LHS = RHS;
+ RHS += 1;
+ }
+
+ // { v = x; x = x binop expr; }
+#pragma acc atomic capture
+ {
+ LHS = RHS;
+ RHS = RHS / 1;
+ }
+#pragma acc atomic capture
+ {
+ LHS = RHS;
+ // expected-error at -3{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{left hand side of assignment operation('RHS') must match one side of the sub-operation on the right hand side('LHS' and '1')}}
+ RHS = LHS ^ 1;
+ }
+#pragma acc atomic capture
+ {
+ LHS = RHS;
+ // expected-error at -3{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{left hand side of assignment operation('LHS') must match one side of the sub-operation on the right hand side('RHS' and '1')}}
+ LHS = RHS << 1;
+ }
+ // { v = x; x = expr binop x; }
+#pragma acc atomic capture
+ {
+ LHS = RHS;
+ RHS = 1 / RHS;
+ }
+#pragma acc atomic capture
+ {
+ LHS = RHS;
+ // expected-error at -3{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{left hand side of assignment operation('RHS') must match one side of the sub-operation on the right hand side('1' and 'LHS')}}
+ RHS = 1 ^ LHS;
+ }
+#pragma acc atomic capture
+ {
+ LHS = RHS;
+ // expected-error at -3{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{left hand side of assignment operation('LHS') must match one side of the sub-operation on the right hand side('1' and 'RHS')}}
+ LHS = 1 << RHS;
+ }
+
+ // { v = x; x = expr; }
+#pragma acc atomic capture
+ {
+ LHS = RHS;
+ RHS = 1;
+ }
+#pragma acc atomic capture
+ {
+ LHS = RHS;
+ // expected-error at -3{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{variable on left hand side of assignment('LHS') must match variable used on right hand side of assignment('RHS') from the first statement}}
+ LHS = 1;
+ }
+
+ // { v = x; x++; }
+ // { v = x; ++x; }
+ // { v = x; x--; }
+ // { v = x; --x; }
+#pragma acc atomic capture
+ {
+ LHS = RHS;
+ RHS++;
+ }
+#pragma acc atomic capture
+ {
+ LHS = RHS;
+ RHS--;
+ }
+#pragma acc atomic capture
+ {
+ LHS = RHS;
+ ++RHS;
+ }
+#pragma acc atomic capture
+ {
+ LHS = RHS;
+ --RHS;
+ }
+#pragma acc atomic capture
+ {
+ LHS = RHS;
+ // expected-error at -3{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{unary operator not supported, only increment and decrement operations permitted}}
+ -RHS;
+ }
+#pragma acc atomic capture
+ {
+ LHS = RHS;
+ // expected-error at -3{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{variable in unary expression('LHS') must match variable used on right hand side of assignment('RHS') from the first statement}}
+ LHS++;
+ }
+}
+
+template<typename T>
+void AtomicCaptureTemplateCompound2(T LHS, T RHS) {
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{expected assignment, compound assignment, increment, or decrement expression}}
+#pragma acc atomic capture
+ {
+ }
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +4{{expected assignment, compound assignment, increment, or decrement expression}}
+#pragma acc atomic capture
+ {
+ LHS = RHS;
+ }
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +3{{'atomic capture' with a compound statement only supports two statements}}
+#pragma acc atomic capture
+ {
+ LHS = RHS; RHS += 1; LHS=RHS;
+ }
+
+
+#pragma acc atomic capture
+ {
+ // expected-error at -2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{operand to increment expression must be of scalar type (was 'Struct')}}
+ LHS++;
+ RHS = LHS;
+ }
+
+#pragma acc atomic capture
+ {
+ // expected-error at -2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{operand to increment expression must be of scalar type (was 'Struct')}}
+ ++LHS;
+ RHS = LHS;
+ }
+
+#pragma acc atomic capture
+ {
+ // expected-error at -2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{operand to decrement expression must be of scalar type (was 'Struct')}}
+ --LHS;
+ RHS = LHS;
+ }
+
+
+#pragma acc atomic capture
+ {
+ LHS--;
+ // expected-error at -3{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{variable on right hand side of assignment('RHS') must match variable used in unary expression('LHS') from the first statement}}
+ LHS = RHS;
+ }
+
+#pragma acc atomic capture
+ {
+ // expected-error at -2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{unary operator not supported, only increment and decrement operations permitted}}
+ -LHS;
+ RHS = LHS;
+ }
+
+#pragma acc atomic capture
+ {
+ --LHS;
+ // expected-error at -3{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{expected assignment expression}}
+ RHS += LHS;
+ }
+
+ // { x binop = expr; v = x; }
+#pragma acc atomic capture
+ {
+ // expected-error at -2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{left operand to compound assignment expression must be of scalar type (was 'Struct')}}
+ LHS += 1;
+ RHS = LHS;
+ }
+#pragma acc atomic capture
+ {
+ LHS *= 1;
+ // expected-error at -3{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{variable on right hand side of assignment('RHS') must match variable used on left hand side of compound assignment('LHS') from the first statement}}
+ LHS = RHS;
+ }
+#pragma acc atomic capture
+ {
+ LHS /= 1;
+ // expected-error at -3{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{expected assignment expression}}
+ RHS += LHS;
+ }
+
+ // { x = x binop expr; v = x; }
+#pragma acc atomic capture
+ {
+ // expected-error at -2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{left operand to assignment expression must be of scalar type (was 'Struct')}}
+ LHS = LHS + 1;
+ RHS = LHS;
+ }
+
+#pragma acc atomic capture
+ {
+ // expected-error at -2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{left hand side of assignment operation('LHS') must match one side of the sub-operation on the right hand side('RHS' and '1')}}
+ LHS = RHS - 1;
+ RHS = LHS;
+ }
+#pragma acc atomic capture
+ {
+ LHS = LHS * 1;
+ // expected-error at -3{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{variable on right hand side of assignment('RHS') must match variable used on left hand side of assignment('LHS') from the first statement}}
+ RHS = RHS;
+ }
+#pragma acc atomic capture
+ {
+ LHS = LHS / 1;
+ // expected-error at -3{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{expected assignment expression}}
+ RHS += LHS;
+ }
+
+ // { x = expr binop x; v = x; }
+#pragma acc atomic capture
+ {
+ // expected-error at -2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{left operand to assignment expression must be of scalar type (was 'Struct')}}
+ LHS = 1 ^ LHS;
+ RHS = LHS;
+ }
+
+#pragma acc atomic capture
+ {
+ // expected-error at -2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{left hand side of assignment operation('LHS') must match one side of the sub-operation on the right hand side('1' and 'RHS')}}
+ LHS = 1 & RHS;
+ RHS = LHS;
+ }
+#pragma acc atomic capture
+ {
+ LHS = LHS | 1;
+ // expected-error at -3{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{variable on right hand side of assignment('RHS') must match variable used on left hand side of assignment('LHS') from the first statement}}
+ RHS = RHS;
+ }
+#pragma acc atomic capture
+ {
+ LHS = LHS << 1;
+ // expected-error at -3{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{expected assignment expression}}
+ RHS += LHS;
+ }
+
+ // { v = x; x binop = expr; }
+#pragma acc atomic capture
+ {
+ // expected-error at -2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{left operand to assignment expression must be of scalar type (was 'Struct')}}
+ LHS = RHS;
+ RHS += 1;
+ }
+
+ // { v = x; x = x binop expr; }
+#pragma acc atomic capture
+ {
+ // expected-error at -2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{left operand to assignment expression must be of scalar type (was 'Struct')}}
+ LHS = RHS;
+ RHS = RHS / 1;
+ }
+#pragma acc atomic capture
+ {
+ LHS = RHS;
+ // expected-error at -3{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{left hand side of assignment operation('RHS') must match one side of the sub-operation on the right hand side('LHS' and '1')}}
+ RHS = LHS ^ 1;
+ }
+#pragma acc atomic capture
+ {
+ LHS = RHS;
+ // expected-error at -3{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{left hand side of assignment operation('LHS') must match one side of the sub-operation on the right hand side('RHS' and '1')}}
+ LHS = RHS << 1;
+ }
+ // { v = x; x = expr binop x; }
+#pragma acc atomic capture
+ {
+ // expected-error at -2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{left operand to assignment expression must be of scalar type (was 'Struct')}}
+ LHS = RHS;
+ RHS = 1 / RHS;
+ }
+#pragma acc atomic capture
+ {
+ LHS = RHS;
+ // expected-error at -3{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{left hand side of assignment operation('RHS') must match one side of the sub-operation on the right hand side('1' and 'LHS')}}
+ RHS = 1 ^ LHS;
+ }
+#pragma acc atomic capture
+ {
+ LHS = RHS;
+ // expected-error at -3{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{left hand side of assignment operation('LHS') must match one side of the sub-operation on the right hand side('1' and 'RHS')}}
+ LHS = 1 << RHS;
+ }
+
+ // { v = x; x = expr; }
+#pragma acc atomic capture
+ {
+ // expected-error at -2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{left operand to assignment expression must be of scalar type (was 'Struct')}}
+ LHS = RHS;
+ RHS = 1;
+ }
+#pragma acc atomic capture
+ {
+ LHS = RHS;
+ // expected-error at -3{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{variable on left hand side of assignment('LHS') must match variable used on right hand side of assignment('RHS') from the first statement}}
+ LHS = 1;
+ }
+
+ // { v = x; x++; }
+ // { v = x; ++x; }
+ // { v = x; x--; }
+ // { v = x; --x; }
+#pragma acc atomic capture
+ {
+ // expected-error at -2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{left operand to assignment expression must be of scalar type (was 'Struct')}}
+ LHS = RHS;
+ RHS++;
+ }
+#pragma acc atomic capture
+ {
+ // expected-error at -2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{left operand to assignment expression must be of scalar type (was 'Struct')}}
+ LHS = RHS;
+ RHS--;
+ }
+#pragma acc atomic capture
+ {
+ // expected-error at -2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{left operand to assignment expression must be of scalar type (was 'Struct')}}
+ LHS = RHS;
+ ++RHS;
+ }
+#pragma acc atomic capture
+ {
+ // expected-error at -2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{left operand to assignment expression must be of scalar type (was 'Struct')}}
+ LHS = RHS;
+ --RHS;
+ }
+#pragma acc atomic capture
+ {
+ LHS = RHS;
+ // expected-error at -3{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{unary operator not supported, only increment and decrement operations permitted}}
+ -RHS;
+ }
+#pragma acc atomic capture
+ {
+ LHS = RHS;
+ // expected-error at -3{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{variable in unary expression('LHS') must match variable used on right hand side of assignment('RHS') from the first statement}}
+ LHS++;
+ }
+}
+void AtomicCaptureCompound(int LHS, int RHS) {
+ AtomicCaptureTemplateCompound(1, 2);
+ AtomicCaptureTemplateCompound2(S1, S2); //expected-note{{in instantiation of function template specialization}}
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +2{{expected assignment, compound assignment, increment, or decrement expression}}
+#pragma acc atomic capture
+ {
+ }
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +4{{expected assignment, compound assignment, increment, or decrement expression}}
+#pragma acc atomic capture
+ {
+ LHS = RHS;
+ }
+
+ // expected-error at +2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +3{{'atomic capture' with a compound statement only supports two statements}}
+#pragma acc atomic capture
+ {
+ LHS = RHS; RHS += 1; LHS=RHS;
+ }
+
+
+#pragma acc atomic capture
+ {
+ LHS++;
+ RHS = LHS;
+ }
+#pragma acc atomic capture
+ {
+ // expected-error at -2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{operand to increment expression must be of scalar type (was 'Struct')}}
+ S1++;
+ S2= S1;
+ }
+
+#pragma acc atomic capture
+ {
+ ++LHS;
+ RHS = LHS;
+ }
+
+#pragma acc atomic capture
+ {
+ --LHS;
+ RHS = LHS;
+ }
+
+
+#pragma acc atomic capture
+ {
+ LHS--;
+ // expected-error at -3{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{variable on right hand side of assignment('RHS') must match variable used in unary expression('LHS') from the first statement}}
+ LHS = RHS;
+ }
+
+#pragma acc atomic capture
+ {
+ // expected-error at -2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{unary operator not supported, only increment and decrement operations permitted}}
+ -LHS;
+ RHS = LHS;
+ }
+
+#pragma acc atomic capture
+ {
+ --LHS;
+ // expected-error at -3{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{expected assignment expression}}
+ RHS += LHS;
+ }
+
+ // { x binop = expr; v = x; }
+#pragma acc atomic capture
+ {
+ LHS += 1;
+ RHS = LHS;
+ }
+#pragma acc atomic capture
+ {
+ // expected-error at -2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{left operand to compound assignment expression must be of scalar type (was 'Struct')}}
+ S1 += 1;
+ S2= S1;
+ }
+#pragma acc atomic capture
+ {
+ LHS *= 1;
+ // expected-error at -3{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{variable on right hand side of assignment('RHS') must match variable used on left hand side of compound assignment('LHS') from the first statement}}
+ LHS = RHS;
+ }
+#pragma acc atomic capture
+ {
+ LHS /= 1;
+ // expected-error at -3{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{expected assignment expression}}
+ RHS += LHS;
+ }
+
+ // { x = x binop expr; v = x; }
+#pragma acc atomic capture
+ {
+ LHS = LHS + 1;
+ RHS = LHS;
+ }
+#pragma acc atomic capture
+ {
+ // expected-error at -2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{left operand to assignment expression must be of scalar type (was 'Struct')}}
+ S1 = S1 + 1;
+ S2= S1;
+ }
+
+#pragma acc atomic capture
+ {
+ // expected-error at -2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{left hand side of assignment operation('LHS') must match one side of the sub-operation on the right hand side('RHS' and '1')}}
+ LHS = RHS - 1;
+ RHS = LHS;
+ }
+#pragma acc atomic capture
+ {
+ LHS = LHS * 1;
+ // expected-error at -3{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{variable on right hand side of assignment('RHS') must match variable used on left hand side of assignment('LHS') from the first statement}}
+ RHS = RHS;
+ }
+#pragma acc atomic capture
+ {
+ LHS = LHS / 1;
+ // expected-error at -3{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{expected assignment expression}}
+ RHS += LHS;
+ }
+
+ // { x = expr binop x; v = x; }
+#pragma acc atomic capture
+ {
+ LHS = 1 ^ LHS;
+ RHS = LHS;
+ }
+#pragma acc atomic capture
+ {
+ // expected-error at -2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{left operand to assignment expression must be of scalar type (was 'Struct')}}
+ S1 = 1 ^ S1;
+ S2 = S1;
+ }
+
+#pragma acc atomic capture
+ {
+ // expected-error at -2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{left hand side of assignment operation('LHS') must match one side of the sub-operation on the right hand side('1' and 'RHS')}}
+ LHS = 1 & RHS;
+ RHS = LHS;
+ }
+#pragma acc atomic capture
+ {
+ LHS = LHS | 1;
+ // expected-error at -3{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{variable on right hand side of assignment('RHS') must match variable used on left hand side of assignment('LHS') from the first statement}}
+ RHS = RHS;
+ }
+#pragma acc atomic capture
+ {
+ LHS = LHS << 1;
+ // expected-error at -3{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{expected assignment expression}}
+ RHS += LHS;
+ }
+
+ // { v = x; x binop = expr; }
+#pragma acc atomic capture
+ {
+ LHS = RHS;
+ RHS += 1;
+ }
+
+#pragma acc atomic capture
+ {
+ // expected-error at -2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{left operand to assignment expression must be of scalar type (was 'Struct')}}
+ S1 = S2;
+ S2 += 1;
+ }
+
+ // { v = x; x = x binop expr; }
+#pragma acc atomic capture
+ {
+ LHS = RHS;
+ RHS = RHS / 1;
+ }
+#pragma acc atomic capture
+ {
+ LHS = RHS;
+ // expected-error at -3{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{left hand side of assignment operation('RHS') must match one side of the sub-operation on the right hand side('LHS' and '1')}}
+ RHS = LHS ^ 1;
+ }
+#pragma acc atomic capture
+ {
+ LHS = RHS;
+ // expected-error at -3{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{left hand side of assignment operation('LHS') must match one side of the sub-operation on the right hand side('RHS' and '1')}}
+ LHS = RHS << 1;
+ }
+ // { v = x; x = expr binop x; }
+#pragma acc atomic capture
+ {
+ LHS = RHS;
+ RHS = 1 / RHS;
+ }
+#pragma acc atomic capture
+ {
+ LHS = RHS;
+ // expected-error at -3{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{left hand side of assignment operation('RHS') must match one side of the sub-operation on the right hand side('1' and 'LHS')}}
+ RHS = 1 ^ LHS;
+ }
+#pragma acc atomic capture
+ {
+ LHS = RHS;
+ // expected-error at -3{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{left hand side of assignment operation('LHS') must match one side of the sub-operation on the right hand side('1' and 'RHS')}}
+ LHS = 1 << RHS;
+ }
+
+ // { v = x; x = expr; }
+#pragma acc atomic capture
+ {
+ LHS = RHS;
+ RHS = 1;
+ }
+#pragma acc atomic capture
+ {
+ LHS = RHS;
+ // expected-error at -3{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{variable on left hand side of assignment('LHS') must match variable used on right hand side of assignment('RHS') from the first statement}}
+ LHS = 1;
+ }
+
+ // { v = x; x++; }
+ // { v = x; ++x; }
+ // { v = x; x--; }
+ // { v = x; --x; }
+#pragma acc atomic capture
+ {
+ LHS = RHS;
+ RHS++;
+ }
+#pragma acc atomic capture
+ {
+ LHS = RHS;
+ RHS--;
+ }
+#pragma acc atomic capture
+ {
+ LHS = RHS;
+ ++RHS;
+ }
+#pragma acc atomic capture
+ {
+ LHS = RHS;
+ --RHS;
+ }
+#pragma acc atomic capture
+ {
+ // expected-error at -2{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{left operand to assignment expression must be of scalar type (was 'Struct')}}
+ S1= S2;
+ --S2;
+ }
+#pragma acc atomic capture
+ {
+ LHS = RHS;
+ // expected-error at -3{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{unary operator not supported, only increment and decrement operations permitted}}
+ -RHS;
+ }
+#pragma acc atomic capture
+ {
+ LHS = RHS;
+ // expected-error at -3{{statement associated with OpenACC 'atomic capture' directive is invalid}}
+ // expected-note at +1{{variable in unary expression('LHS') must match variable used on right hand side of assignment('RHS') from the first statement}}
+ LHS++;
+ }
+}
diff --git a/clang/tools/libclang/CIndex.cpp b/clang/tools/libclang/CIndex.cpp
index bf7fdeec0cc51b0..697cc4776839dcd 100644
--- a/clang/tools/libclang/CIndex.cpp
+++ b/clang/tools/libclang/CIndex.cpp
@@ -2193,6 +2193,8 @@ class EnqueueVisitor : public ConstStmtVisitor<EnqueueVisitor, void>,
void VisitOpenACCInitConstruct(const OpenACCInitConstruct *D);
void VisitOpenACCShutdownConstruct(const OpenACCShutdownConstruct *D);
void VisitOpenACCSetConstruct(const OpenACCSetConstruct *D);
+ void VisitOpenACCUpdateConstruct(const OpenACCUpdateConstruct *D);
+ void VisitOpenACCAtomicConstruct(const OpenACCAtomicConstruct *D);
void VisitOMPExecutableDirective(const OMPExecutableDirective *D);
void VisitOMPLoopBasedDirective(const OMPLoopBasedDirective *D);
void VisitOMPLoopDirective(const OMPLoopDirective *D);
@@ -3682,6 +3684,18 @@ void EnqueueVisitor::VisitOpenACCSetConstruct(const OpenACCSetConstruct *C) {
EnqueueChildren(Clause);
}
+void EnqueueVisitor::VisitOpenACCUpdateConstruct(
+ const OpenACCUpdateConstruct *C) {
+ EnqueueChildren(C);
+ for (auto *Clause : C->clauses())
+ EnqueueChildren(Clause);
+}
+
+void EnqueueVisitor::VisitOpenACCAtomicConstruct(
+ const OpenACCAtomicConstruct *C) {
+ EnqueueChildren(C);
+}
+
void EnqueueVisitor::VisitAnnotateAttr(const AnnotateAttr *A) {
EnqueueChildren(A);
}
@@ -6454,6 +6468,8 @@ CXString clang_getCursorKindSpelling(enum CXCursorKind Kind) {
return cxstring::createRef("OpenACCSetConstruct");
case CXCursor_OpenACCUpdateConstruct:
return cxstring::createRef("OpenACCUpdateConstruct");
+ case CXCursor_OpenACCAtomicConstruct:
+ return cxstring::createRef("OpenACCAtomicConstruct");
}
llvm_unreachable("Unhandled CXCursorKind");
diff --git a/clang/tools/libclang/CXCursor.cpp b/clang/tools/libclang/CXCursor.cpp
index 60c740311e940b6..127f22bc5bdca06 100644
--- a/clang/tools/libclang/CXCursor.cpp
+++ b/clang/tools/libclang/CXCursor.cpp
@@ -920,6 +920,9 @@ CXCursor cxcursor::MakeCXCursor(const Stmt *S, const Decl *Parent,
case Stmt::OpenACCUpdateConstructClass:
K = CXCursor_OpenACCUpdateConstruct;
break;
+ case Stmt::OpenACCAtomicConstructClass:
+ K = CXCursor_OpenACCAtomicConstruct;
+ break;
case Stmt::OMPTargetParallelGenericLoopDirectiveClass:
K = CXCursor_OMPTargetParallelGenericLoopDirective;
break;
More information about the cfe-commits
mailing list