[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