diff --git a/clang/include/clang/AST/StmtOpenACC.h b/clang/include/clang/AST/StmtOpenACC.h
index c2c74f5cf195..39c4c8184491 100644
--- a/clang/include/clang/AST/StmtOpenACC.h
+++ b/clang/include/clang/AST/StmtOpenACC.h
@@ -829,24 +829,42 @@ public:
 
 // This class represents the 'atomic' construct, which has an associated
 // statement, but no clauses.
-class OpenACCAtomicConstruct final : public OpenACCAssociatedStmtConstruct {
+class OpenACCAtomicConstruct final
+    : public OpenACCAssociatedStmtConstruct,
+      private llvm::TrailingObjects<OpenACCAtomicConstruct,
+                                    const OpenACCClause *> {
 
   friend class ASTStmtReader;
+  friend TrailingObjects;
   OpenACCAtomicKind AtomicKind = OpenACCAtomicKind::None;
 
-  OpenACCAtomicConstruct(EmptyShell)
+  OpenACCAtomicConstruct(unsigned NumClauses)
       : OpenACCAssociatedStmtConstruct(
             OpenACCAtomicConstructClass, OpenACCDirectiveKind::Atomic,
             SourceLocation{}, SourceLocation{}, SourceLocation{},
-            /*AssociatedStmt=*/nullptr) {}
+            /*AssociatedStmt=*/nullptr) {
+    std::uninitialized_value_construct(
+        getTrailingObjects<const OpenACCClause *>(),
+        getTrailingObjects<const OpenACCClause *>() + NumClauses);
+    setClauseList(MutableArrayRef(getTrailingObjects<const OpenACCClause *>(),
+                                  NumClauses));
+  }
 
   OpenACCAtomicConstruct(SourceLocation Start, SourceLocation DirectiveLoc,
                          OpenACCAtomicKind AtKind, SourceLocation End,
+                         ArrayRef<const OpenACCClause *> Clauses,
                          Stmt *AssociatedStmt)
       : OpenACCAssociatedStmtConstruct(OpenACCAtomicConstructClass,
                                        OpenACCDirectiveKind::Atomic, Start,
                                        DirectiveLoc, End, AssociatedStmt),
-        AtomicKind(AtKind) {}
+        AtomicKind(AtKind) {
+    // Initialize the trailing storage.
+    std::uninitialized_copy(Clauses.begin(), Clauses.end(),
+                            getTrailingObjects<const OpenACCClause *>());
+
+    setClauseList(MutableArrayRef(getTrailingObjects<const OpenACCClause *>(),
+                                  Clauses.size()));
+  }
 
   void setAssociatedStmt(Stmt *S) {
     OpenACCAssociatedStmtConstruct::setAssociatedStmt(S);
@@ -857,10 +875,12 @@ public:
     return T->getStmtClass() == OpenACCAtomicConstructClass;
   }
 
-  static OpenACCAtomicConstruct *CreateEmpty(const ASTContext &C);
+  static OpenACCAtomicConstruct *CreateEmpty(const ASTContext &C,
+                                             unsigned NumClauses);
   static OpenACCAtomicConstruct *
   Create(const ASTContext &C, SourceLocation Start, SourceLocation DirectiveLoc,
-         OpenACCAtomicKind AtKind, SourceLocation End, Stmt *AssociatedStmt);
+         OpenACCAtomicKind AtKind, SourceLocation End,
+         ArrayRef<const OpenACCClause *> Clauses, Stmt *AssociatedStmt);
 
   OpenACCAtomicKind getAtomicKind() const { return AtomicKind; }
   const Stmt *getAssociatedStmt() const {
diff --git a/clang/lib/AST/StmtOpenACC.cpp b/clang/lib/AST/StmtOpenACC.cpp
index 8a86074fe68a..c45eca92dc87 100644
--- a/clang/lib/AST/StmtOpenACC.cpp
+++ b/clang/lib/AST/StmtOpenACC.cpp
@@ -307,20 +307,26 @@ OpenACCUpdateConstruct::Create(const ASTContext &C, SourceLocation Start,
 }
 
 OpenACCAtomicConstruct *
-OpenACCAtomicConstruct::CreateEmpty(const ASTContext &C) {
-  void *Mem = C.Allocate(sizeof(OpenACCAtomicConstruct));
-  auto *Inst = new (Mem) OpenACCAtomicConstruct(EmptyShell{});
+OpenACCAtomicConstruct::CreateEmpty(const ASTContext &C, unsigned NumClauses) {
+  void *Mem = C.Allocate(
+      OpenACCAtomicConstruct::totalSizeToAlloc<const OpenACCClause *>(
+          NumClauses));
+  auto *Inst = new (Mem) OpenACCAtomicConstruct(NumClauses);
   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);
+    OpenACCAtomicKind AtKind, SourceLocation End,
+    ArrayRef<const OpenACCClause *> Clauses, Stmt *AssociatedStmt) {
+  void *Mem = C.Allocate(
+      OpenACCAtomicConstruct::totalSizeToAlloc<const OpenACCClause *>(
+          Clauses.size()));
+  auto *Inst = new (Mem) OpenACCAtomicConstruct(Start, DirectiveLoc, AtKind,
+                                                End, Clauses, AssociatedStmt);
   return Inst;
 }
+
 OpenACCCacheConstruct *OpenACCCacheConstruct::CreateEmpty(const ASTContext &C,
                                                           unsigned NumVars) {
   void *Mem =
diff --git a/clang/lib/AST/StmtPrinter.cpp b/clang/lib/AST/StmtPrinter.cpp
index fe9784a15b76..dbe2432d5c79 100644
--- a/clang/lib/AST/StmtPrinter.cpp
+++ b/clang/lib/AST/StmtPrinter.cpp
@@ -1258,6 +1258,7 @@ void StmtPrinter::VisitOpenACCAtomicConstruct(OpenACCAtomicConstruct *S) {
   if (S->getAtomicKind() != OpenACCAtomicKind::None)
     OS << " " << S->getAtomicKind();
 
+  PrintOpenACCClauseList(S);
   OS << '\n';
   PrintStmt(S->getAssociatedStmt());
 }
diff --git a/clang/lib/AST/StmtProfile.cpp b/clang/lib/AST/StmtProfile.cpp
index 96852858b3b9..83d54da9be7e 100644
--- a/clang/lib/AST/StmtProfile.cpp
+++ b/clang/lib/AST/StmtProfile.cpp
@@ -2839,6 +2839,8 @@ void StmtProfiler::VisitOpenACCUpdateConstruct(
 void StmtProfiler::VisitOpenACCAtomicConstruct(
     const OpenACCAtomicConstruct *S) {
   VisitStmt(S);
+  OpenACCClauseProfiler P{*this};
+  P.VisitOpenACCClauseList(S->clauses());
 }
 
 void StmtProfiler::VisitHLSLOutArgExpr(const HLSLOutArgExpr *S) {
diff --git a/clang/lib/Sema/SemaOpenACC.cpp b/clang/lib/Sema/SemaOpenACC.cpp
index fea334e9e80f..231e73fdbc9e 100644
--- a/clang/lib/Sema/SemaOpenACC.cpp
+++ b/clang/lib/Sema/SemaOpenACC.cpp
@@ -1832,9 +1832,8 @@ StmtResult SemaOpenACC::ActOnEndStmtDirective(
                                           EndLoc, Clauses);
   }
   case OpenACCDirectiveKind::Atomic: {
-    assert(Clauses.empty() && "Atomic doesn't allow clauses");
     return OpenACCAtomicConstruct::Create(
-        getASTContext(), StartLoc, DirLoc, AtomicKind, EndLoc,
+        getASTContext(), StartLoc, DirLoc, AtomicKind, EndLoc, Clauses,
         AssocStmt.isUsable() ? AssocStmt.get() : nullptr);
   }
   case OpenACCDirectiveKind::Cache: {
diff --git a/clang/lib/Sema/SemaOpenACCClause.cpp b/clang/lib/Sema/SemaOpenACCClause.cpp
index 78e250b75950..a98b6712014c 100644
--- a/clang/lib/Sema/SemaOpenACCClause.cpp
+++ b/clang/lib/Sema/SemaOpenACCClause.cpp
@@ -55,6 +55,8 @@ bool doesClauseApplyToDirective(OpenACCDirectiveKind DirectiveKind,
     case OpenACCDirectiveKind::ParallelLoop:
     case OpenACCDirectiveKind::SerialLoop:
     case OpenACCDirectiveKind::KernelsLoop:
+      // OpenACC 3.4(prerelease) PR #511 adds 'if' to atomic.
+    case OpenACCDirectiveKind::Atomic:
       return true;
     default:
       return false;
diff --git a/clang/lib/Sema/TreeTransform.h b/clang/lib/Sema/TreeTransform.h
index 916b8e2735cd..524e73242a2d 100644
--- a/clang/lib/Sema/TreeTransform.h
+++ b/clang/lib/Sema/TreeTransform.h
@@ -4222,10 +4222,11 @@ public:
                                            SourceLocation DirLoc,
                                            OpenACCAtomicKind AtKind,
                                            SourceLocation EndLoc,
+                                           ArrayRef<OpenACCClause *> Clauses,
                                            StmtResult AssociatedStmt) {
     return getSema().OpenACC().ActOnEndStmtDirective(
         OpenACCDirectiveKind::Atomic, BeginLoc, DirLoc, SourceLocation{},
-        SourceLocation{}, {}, AtKind, SourceLocation{}, EndLoc, {},
+        SourceLocation{}, {}, AtKind, SourceLocation{}, EndLoc, Clauses,
         AssociatedStmt);
   }
 
@@ -12744,6 +12745,10 @@ StmtResult TreeTransform<Derived>::TransformOpenACCAtomicConstruct(
     OpenACCAtomicConstruct *C) {
   getSema().OpenACC().ActOnConstruct(C->getDirectiveKind(), C->getBeginLoc());
 
+  llvm::SmallVector<OpenACCClause *> TransformedClauses =
+      getDerived().TransformOpenACCClauseList(C->getDirectiveKind(),
+                                              C->clauses());
+
   if (getSema().OpenACC().ActOnStartStmtDirective(C->getDirectiveKind(),
                                                   C->getBeginLoc(), {}))
     return StmtError();
@@ -12759,7 +12764,7 @@ StmtResult TreeTransform<Derived>::TransformOpenACCAtomicConstruct(
 
   return getDerived().RebuildOpenACCAtomicConstruct(
       C->getBeginLoc(), C->getDirectiveLoc(), C->getAtomicKind(),
-      C->getEndLoc(), AssocStmt);
+      C->getEndLoc(), TransformedClauses, AssocStmt);
 }
 
 template <typename Derived>
diff --git a/clang/lib/Serialization/ASTReaderStmt.cpp b/clang/lib/Serialization/ASTReaderStmt.cpp
index 8dceca6ff3db..7c7abcb2d49d 100644
--- a/clang/lib/Serialization/ASTReaderStmt.cpp
+++ b/clang/lib/Serialization/ASTReaderStmt.cpp
@@ -2940,9 +2940,7 @@ void ASTStmtReader::VisitOpenACCCacheConstruct(OpenACCCacheConstruct *S) {
 
 void ASTStmtReader::VisitOpenACCAtomicConstruct(OpenACCAtomicConstruct *S) {
   VisitStmt(S);
-  S->Kind = Record.readEnum<OpenACCDirectiveKind>();
-  S->Range = Record.readSourceRange();
-  S->DirectiveLoc = Record.readSourceLocation();
+  VisitOpenACCConstructStmt(S);
   S->AtomicKind = Record.readEnum<OpenACCAtomicKind>();
   S->setAssociatedStmt(Record.readSubStmt());
 }
@@ -4490,7 +4488,8 @@ Stmt *ASTReader::ReadStmtFromStream(ModuleFile &F) {
       break;
     }
     case STMT_OPENACC_ATOMIC_CONSTRUCT: {
-      S = OpenACCAtomicConstruct::CreateEmpty(Context);
+      unsigned NumClauses = Record[ASTStmtReader::NumStmtFields];
+      S = OpenACCAtomicConstruct::CreateEmpty(Context, NumClauses);
       break;
     }
     case EXPR_REQUIRES: {
diff --git a/clang/lib/Serialization/ASTWriterStmt.cpp b/clang/lib/Serialization/ASTWriterStmt.cpp
index 037eb13b1dc5..0860704368f3 100644
--- a/clang/lib/Serialization/ASTWriterStmt.cpp
+++ b/clang/lib/Serialization/ASTWriterStmt.cpp
@@ -3014,9 +3014,7 @@ void ASTStmtWriter::VisitOpenACCWaitConstruct(OpenACCWaitConstruct *S) {
 
 void ASTStmtWriter::VisitOpenACCAtomicConstruct(OpenACCAtomicConstruct *S) {
   VisitStmt(S);
-  Record.writeEnum(S->Kind);
-  Record.AddSourceRange(S->Range);
-  Record.AddSourceLocation(S->DirectiveLoc);
+  VisitOpenACCConstructStmt(S);
   Record.writeEnum(S->getAtomicKind());
   Record.AddStmt(S->getAssociatedStmt());
 
diff --git a/clang/test/AST/ast-print-openacc-atomic-construct.cpp b/clang/test/AST/ast-print-openacc-atomic-construct.cpp
index 572f2ea4842d..50e988be603d 100644
--- a/clang/test/AST/ast-print-openacc-atomic-construct.cpp
+++ b/clang/test/AST/ast-print-openacc-atomic-construct.cpp
@@ -31,3 +31,35 @@ void foo(int v, int x) {
   { x--; v = x; }
 
 }
+
+void foo2(int v, int x) {
+// CHECK: #pragma acc atomic read if(v)
+// CHECK-NEXT:   v = x;
+#pragma acc atomic read if (v)
+  v = x;
+// CHECK-NEXT: pragma acc atomic write if(x)
+// CHECK-NEXT:  v = x + 1;
+#pragma acc atomic write if (x)
+  v = x + 1;
+// CHECK-NEXT: pragma acc atomic update if(true)
+// CHECK-NEXT:  x++;
+#pragma acc atomic update if (true)
+  x++;
+// CHECK-NEXT: pragma acc atomic  if(false)
+// CHECK-NEXT:  x--;
+#pragma acc atomic if (false)
+  x--;
+// CHECK-NEXT: pragma acc atomic capture if(v < x)
+// CHECK-NEXT:  v = x++;
+#pragma acc atomic capture if (v < x)
+  v = x++;
+
+// CHECK-NEXT: #pragma acc atomic capture if(x > v)
+// CHECK-NEXT: { 
+// CHECK-NEXT: x--;
+// CHECK-NEXT: v = x;
+// CHECK-NEXT: }
+#pragma acc atomic capture if (x > v)
+  { x--; v = x; }
+
+}
diff --git a/clang/test/ParserOpenACC/parse-clauses.c b/clang/test/ParserOpenACC/parse-clauses.c
index 100dc07a9fd6..e31b7492dab2 100644
--- a/clang/test/ParserOpenACC/parse-clauses.c
+++ b/clang/test/ParserOpenACC/parse-clauses.c
@@ -1288,3 +1288,27 @@ void BCP1();
 #pragma acc routine seq bind("ReductionClauseParsing")
 
 #pragma acc routine(BCP1) seq bind(unknown_thing)
+
+void AtomicIf() {
+  int i, j;
+  // expected-error@+1{{expected '('}}
+#pragma acc atomic read if
+  i = j;
+#pragma acc atomic read if (0)
+  i = j;
+#pragma acc atomic write if (1)
+  i = j + 1;
+
+#pragma acc atomic update if (i)
+  ++i;
+#pragma acc atomic if (j)
+  ++i;
+
+#pragma acc atomic capture if (0)
+  i = j++;
+#pragma acc atomic capture if (i)
+  {
+    ++j;
+    i = j;
+  }
+}
diff --git a/clang/test/ParserOpenACC/parse-clauses.cpp b/clang/test/ParserOpenACC/parse-clauses.cpp
index 6c8401806094..0ba70f5c6a1f 100644
--- a/clang/test/ParserOpenACC/parse-clauses.cpp
+++ b/clang/test/ParserOpenACC/parse-clauses.cpp
@@ -61,3 +61,27 @@ void use() {
 void another_func();
   // expected-warning@+1{{encoding prefix 'U' on an unevaluated string literal has no effect}}
 #pragma acc routine(another_func) seq bind(U"32 bits")
+
+void AtomicIf() {
+  int i, j;
+  // expected-error@+1{{expected '('}}
+#pragma acc atomic read if
+  i = j;
+#pragma acc atomic read if (true)
+  i = j;
+#pragma acc atomic write if (false)
+  i = j + 1;
+
+#pragma acc atomic update if (i)
+  ++i;
+#pragma acc atomic if (j)
+  ++i;
+
+#pragma acc atomic capture if (true)
+  i = j++;
+#pragma acc atomic capture if (i)
+  {
+    ++j;
+    i = j;
+  }
+}
diff --git a/clang/test/SemaOpenACC/atomic-construct-ast.cpp b/clang/test/SemaOpenACC/atomic-construct-ast.cpp
index bfd7f32884f1..157532951a16 100644
--- a/clang/test/SemaOpenACC/atomic-construct-ast.cpp
+++ b/clang/test/SemaOpenACC/atomic-construct-ast.cpp
@@ -68,6 +68,90 @@ _Pragma("acc atomic read")
 #pragma acc atomic capture
   { x--; v = x; }
 
+// CHECK-NEXT: OpenACCAtomicConstruct{{.*}} atomic read
+// CHECK-NEXT: if clause
+// CHECK-NEXT: BinaryOperator{{.*}}'bool' '<'
+// CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <LValueToRValue>
+// CHECK-NEXT: DeclRefExpr{{.*}}'v' 'int'
+// CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <LValueToRValue>
+// 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 read if (v < x)
+  v = x;
+
+// CHECK-NEXT: OpenACCAtomicConstruct{{.*}} atomic write
+// CHECK-NEXT: if clause
+// CHECK-NEXT: BinaryOperator{{.*}}'bool' '<'
+// CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <LValueToRValue>
+// CHECK-NEXT: DeclRefExpr{{.*}}'v' 'int'
+// CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <LValueToRValue>
+// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'int'
+// 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 if (v < x)")
+  v = x + 1;
+
+// CHECK-NEXT: OpenACCAtomicConstruct{{.*}} atomic update 
+// CHECK-NEXT: if clause
+// CHECK-NEXT: BinaryOperator{{.*}}'bool' '<'
+// CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <LValueToRValue>
+// CHECK-NEXT: DeclRefExpr{{.*}}'v' 'int'
+// CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <LValueToRValue>
+// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'int'
+// CHECK-NEXT: UnaryOperator{{.*}} 'int' postfix '++'
+// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'int'
+#pragma acc atomic update if (v < x)
+  x++;
+
+// CHECK-NEXT: OpenACCAtomicConstruct{{.*}} atomic <none>
+// CHECK-NEXT: if clause
+// CHECK-NEXT: BinaryOperator{{.*}}'bool' '<'
+// CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <LValueToRValue>
+// CHECK-NEXT: DeclRefExpr{{.*}}'v' 'int'
+// CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <LValueToRValue>
+// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'int'
+// CHECK-NEXT: UnaryOperator{{.*}} 'int' postfix '--'
+// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'int'
+#pragma acc atomic if (v < x)
+  x--;
+
+// CHECK-NEXT: OpenACCAtomicConstruct{{.*}} atomic capture
+// CHECK-NEXT: if clause
+// CHECK-NEXT: BinaryOperator{{.*}}'bool' '<'
+// CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <LValueToRValue>
+// CHECK-NEXT: DeclRefExpr{{.*}}'v' 'int'
+// CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <LValueToRValue>
+// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'int'
+// CHECK-NEXT: BinaryOperator{{.*}} 'int' lvalue '='
+// CHECK-NEXT: DeclRefExpr{{.*}}'v' 'int'
+// CHECK-NEXT: UnaryOperator{{.*}} 'int' postfix '++'
+// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'int'
+#pragma acc atomic capture if (v < x)
+  v = x++;
+
+// CHECK-NEXT: OpenACCAtomicConstruct{{.*}} atomic capture
+// CHECK-NEXT: if clause
+// CHECK-NEXT: BinaryOperator{{.*}}'bool' '<'
+// CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <LValueToRValue>
+// CHECK-NEXT: DeclRefExpr{{.*}}'v' 'int'
+// CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <LValueToRValue>
+// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'int'
+// 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 if (v < x)
+  { x--; v = x; }
 }
 
 template<typename T, int I>
@@ -124,6 +208,76 @@ void templ_foo(T v, T x) {
 #pragma acc atomic capture
   { x--; v = x; }
 
+// CHECK-NEXT: OpenACCAtomicConstruct{{.*}} atomic read
+// CHECK-NEXT: if clause
+// CHECK-NEXT: BinaryOperator{{.*}}'<dependent type>' '<'
+// CHECK-NEXT: DeclRefExpr{{.*}}'v' 'T'
+// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'T'
+// CHECK-NEXT: BinaryOperator{{.*}} '<dependent type>' '='
+// CHECK-NEXT: DeclRefExpr{{.*}}'v' 'T'
+// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'T'
+#pragma acc atomic read if (v < x)
+  v = x;
+
+// CHECK-NEXT: OpenACCAtomicConstruct{{.*}} atomic write
+// CHECK-NEXT: if clause
+// CHECK-NEXT: BinaryOperator{{.*}}'<dependent type>' '<'
+// CHECK-NEXT: DeclRefExpr{{.*}}'v' 'T'
+// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'T'
+// 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 if (v < x)")
+  v = x + I;
+
+// CHECK-NEXT: OpenACCAtomicConstruct{{.*}} atomic update 
+// CHECK-NEXT: if clause
+// CHECK-NEXT: BinaryOperator{{.*}}'<dependent type>' '<'
+// CHECK-NEXT: DeclRefExpr{{.*}}'v' 'T'
+// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'T'
+// CHECK-NEXT: UnaryOperator{{.*}} '<dependent type>' postfix '++'
+// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'T'
+#pragma acc atomic update if (v < x)
+  x++;
+
+// CHECK-NEXT: OpenACCAtomicConstruct{{.*}} atomic <none>
+// CHECK-NEXT: if clause
+// CHECK-NEXT: BinaryOperator{{.*}}'<dependent type>' '<'
+// CHECK-NEXT: DeclRefExpr{{.*}}'v' 'T'
+// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'T'
+// CHECK-NEXT: UnaryOperator{{.*}} '<dependent type>' postfix '--'
+// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'T'
+#pragma acc atomic if (v < x)
+  x--;
+
+// CHECK-NEXT: OpenACCAtomicConstruct{{.*}} atomic capture
+// CHECK-NEXT: if clause
+// CHECK-NEXT: BinaryOperator{{.*}}'<dependent type>' '<'
+// CHECK-NEXT: DeclRefExpr{{.*}}'v' 'T'
+// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'T'
+// 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 if (v < x)
+  v = x++;
+
+// CHECK-NEXT: OpenACCAtomicConstruct{{.*}} atomic capture
+// CHECK-NEXT: if clause
+// CHECK-NEXT: BinaryOperator{{.*}}'<dependent type>' '<'
+// CHECK-NEXT: DeclRefExpr{{.*}}'v' 'T'
+// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'T'
+// 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 if (v < x)
+  { x--; v = x; }
+
   // CHECK-NEXT: FunctionDecl{{.*}} templ_foo 'void (int, int)' implicit_instantiation
   // CHECK-NEXT: TemplateArgument type 'int'
   // CHECK-NEXT: BuiltinType{{.*}} 'int'
@@ -170,6 +324,81 @@ void templ_foo(T v, T x) {
 // CHECK-NEXT: DeclRefExpr{{.*}}'v' 'int'
 // CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <LValueToRValue>
 // CHECK-NEXT: DeclRefExpr{{.*}}'x' 'int'
+
+// CHECK-NEXT: OpenACCAtomicConstruct{{.*}} atomic read
+// CHECK-NEXT: if clause
+// CHECK-NEXT: BinaryOperator{{.*}}'bool' '<'
+// CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <LValueToRValue>
+// CHECK-NEXT: DeclRefExpr{{.*}}'v' 'int'
+// CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <LValueToRValue>
+// 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'
+
+// CHECK-NEXT: OpenACCAtomicConstruct{{.*}} atomic write
+// CHECK-NEXT: if clause
+// CHECK-NEXT: BinaryOperator{{.*}}'bool' '<'
+// CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <LValueToRValue>
+// CHECK-NEXT: DeclRefExpr{{.*}}'v' 'int'
+// CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <LValueToRValue>
+// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'int'
+// 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: if clause
+// CHECK-NEXT: BinaryOperator{{.*}}'bool' '<'
+// CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <LValueToRValue>
+// CHECK-NEXT: DeclRefExpr{{.*}}'v' 'int'
+// CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <LValueToRValue>
+// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'int'
+// CHECK-NEXT: UnaryOperator{{.*}} 'int' postfix '++'
+// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'int'
+
+// CHECK-NEXT: OpenACCAtomicConstruct{{.*}} atomic <none>
+// CHECK-NEXT: if clause
+// CHECK-NEXT: BinaryOperator{{.*}}'bool' '<'
+// CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <LValueToRValue>
+// CHECK-NEXT: DeclRefExpr{{.*}}'v' 'int'
+// CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <LValueToRValue>
+// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'int'
+// CHECK-NEXT: UnaryOperator{{.*}} 'int' postfix '--'
+// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'int'
+
+// CHECK-NEXT: OpenACCAtomicConstruct{{.*}} atomic capture
+// CHECK-NEXT: if clause
+// CHECK-NEXT: BinaryOperator{{.*}}'bool' '<'
+// CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <LValueToRValue>
+// CHECK-NEXT: DeclRefExpr{{.*}}'v' 'int'
+// CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <LValueToRValue>
+// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'int'
+// 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: if clause
+// CHECK-NEXT: BinaryOperator{{.*}}'bool' '<'
+// CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <LValueToRValue>
+// CHECK-NEXT: DeclRefExpr{{.*}}'v' 'int'
+// CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <LValueToRValue>
+// CHECK-NEXT: DeclRefExpr{{.*}}'x' 'int'
+// 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() {
diff --git a/clang/test/SemaOpenACC/atomic-construct.cpp b/clang/test/SemaOpenACC/atomic-construct.cpp
index 7357d91d704f..83fbc4bff2ec 100644
--- a/clang/test/SemaOpenACC/atomic-construct.cpp
+++ b/clang/test/SemaOpenACC/atomic-construct.cpp
@@ -47,9 +47,30 @@ Struct &operator-(Struct&, Struct&);
 
 Struct S1, S2;
 
+struct NotCondition{} NC;
+
 template<typename T>
 T &getRValue();
 
+void IfClause(int x, int v) {
+  // expected-error@+1{{OpenACC 'seq' clause is not valid on 'atomic' directive}}
+#pragma acc atomic read seq
+  x = v;
+
+  // expected-error@+1{{expected '('}}
+#pragma acc atomic read if
+  x = v;
+
+  // expected-error@+1{{value of type 'struct NotCondition' is not contextually convertible to 'bool'}}
+#pragma acc atomic read if(NC)
+  x = v;
+
+  // expected-error@+2{{OpenACC 'if' clause cannot appear more than once on a 'atomic' directive}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc atomic read if(x) if (v)
+  x = v;
+}
+
 template<typename T>
 void AtomicReadTemplate(T LHS, T RHS) {
 #pragma acc atomic read
diff --git a/clang/tools/libclang/CIndex.cpp b/clang/tools/libclang/CIndex.cpp
index 197ba2cd6856..6ea6447d1d59 100644
--- a/clang/tools/libclang/CIndex.cpp
+++ b/clang/tools/libclang/CIndex.cpp
@@ -3719,6 +3719,8 @@ void EnqueueVisitor::VisitOpenACCUpdateConstruct(
 void EnqueueVisitor::VisitOpenACCAtomicConstruct(
     const OpenACCAtomicConstruct *C) {
   EnqueueChildren(C);
+  for (auto *Clause : C->clauses())
+    EnqueueChildren(Clause);
 }
 
 void EnqueueVisitor::VisitAnnotateAttr(const AnnotateAttr *A) {