[WIP][Clang][OpenMP] Add the support for compare clause in atomic directive

[tianshilei1992]

Summary

NOTE: Major functionalities have been implemented, but still requires some minor tweak.

NOTE: compare capture combination will be in another patch.

[tianshilei1992]

rebase

[jdoerfert]

Is this still WIP or should it be reviewed? Needs a proper commit message.

clang/lib/CodeGen/CGStmtOpenMP.cpp

5797–5800

This looks like something is missing.

[tianshilei1992]

In D102449#2960209, @jdoerfert wrote:

Is this still WIP or should it be reviewed? Needs a proper commit message.

Not yet.

[ABataev]

clang/include/clang/AST/StmtOpenMP.h

2797–2824

It is worth it to pre-commit these and related changes in a separate NFC patch

2848–2851

Can we pack it into a struct? At least X, V etc. params?

clang/lib/CodeGen/CGStmtOpenMP.cpp

5775–5802

I would think about moving most of the atomic codegen functionality to common runtime/OMPBuilder part to be able to override implementation for different targets. It may help to avoid codegen problems with Nvidia/AMD GPUs. Not directly related to this patch though.

[tianshilei1992]

clang/lib/CodeGen/CGStmtOpenMP.cpp

5775–5802

yeah, actually IIRC the OMPIRBuilder already supports atomic operations. Not sure if it supports the compare clause but I'll take a look. If not, I also think it's a better place to sit them.

[tianshilei1992]

I'm going to move code gen into OMPIRBuilder.

[carlo.bertolli]

Thanks for uploading this.

I tried it locally and it is missing proper parsing +sema of the cond-expr-stmt and cond-update-stmt as indicate in the openmp 5.1 spec's. I am sure you are aware of this and this is still WIP.
Do you plan to add support for that? The regression test is not exposing conditional statements, but update ones, so it is not really testing what it should.

You might also want to split parse+sema support from codegen, if it helps.

Many thanks again!

[tianshilei1992]

Yes, there are a couple of issues of implicit cast to be fixed. Current test is just a placeholder. I copied and pasted from the one for update clause and just replaced all updates with compare. I’m also working on that.

[tianshilei1992]

1. Ignore implicit cast in some cases.
2. Add the test case.

What are still missing:

1. Code generation for floating point types.

Question:

1. What atomic instruction(s) should we use for floating point types? Current I can see we do support floating point types for update clause, but the emitted code doesn't look like atomic operations.
2. When do we have to use IgnoreParenImpCasts?

[tianshilei1992]

Another thing is how we deal with a corner case. Say the OpenMP code is written in the following way:

#pragma omp atomic compare
  x = e < x ? e : x;

That's how OpenMP spec defines the atomic operation. x is always in "else statement" of a conditional statement.

Now we need to lower it to LLVM IR, which is atomicrmw operation. Based on the LLVM IR reference, it only supports the atomic operations that x is in the "then statement". For example: x = x > e ? x : e. See the x here is before :. In order to lower the OpenMP statement, we have to do a transformation. In order to swap e and x, we need to transform it to x = e >= x : x : e, a.k.a. x = x <= e : x : e. However, we don't have an atomic operation for <=. We only have <. So if x != e, the result is good.

The incorrectness happens if x == e. Recall at the OpenMP statement, when x == e, the result should be x = x. But if we look at our lowered LLVM IR, x = x < e : x : e, when x == e, it becomes x = e, which doesn't conform with OpenMP spec.

What should we do here?

[tianshilei1992]

rebase

[carlo.bertolli]

In D102449#2977711, @tianshilei1992 wrote:

Another thing is how we deal with a corner case. Say the OpenMP code is written in the following way:

#pragma omp atomic compare
  x = e < x ? e : x;

That's how OpenMP spec defines the atomic operation. x is always in "else statement" of a conditional statement.

Now we need to lower it to LLVM IR, which is atomicrmw operation. Based on the LLVM IR reference, it only supports the atomic operations that x is in the "then statement". For example: x = x > e ? x : e. See the x here is before :. In order to lower the OpenMP statement, we have to do a transformation. In order to swap e and x, we need to transform it to x = e >= x : x : e, a.k.a. x = x <= e : x : e. However, we don't have an atomic operation for <=. We only have <. So if x != e, the result is good.

The incorrectness happens if x == e. Recall at the OpenMP statement, when x == e, the result should be x = x. But if we look at our lowered LLVM IR, x = x < e : x : e, when x == e, it becomes x = e, which doesn't conform with OpenMP spec.

What should we do here?

There seems to be an imbalance in the OpenMP specs, where I would have written something like the followin (page 268 or 5.1 spec's):

x = expr ordop x ? expr : x;
x = x ordop expr ? x : expr;

I tried thinking about a compact solution to your problem, including trying to take advantage of the fact that when x == e we don't need to actually do the assignment (page 272 line 17 of 5.1 spec's and see below), but I was unsuccessful.
Did you consider special casing this transforming it into the equivalent:
#pragma omp critical
{ if (x > e) x = e; }
?

Remember:

For forms where statement is cond-expr-stmt, if the result of the condition implies that the value of x
does not change then the update may not occur.

The right thing to do would be to fix the spec's, but that can't happen until 5.2.

[carlo.bertolli]

This is already a lot of code with parse+sema. I wonder if we should split the patch into two, i.e. 1. parse+sema; 2. code gen? @ABataev ?
It should simplify maintenance of the patch and allow time to extend the OpenMP IR builder.

[ABataev]

In D102449#3137713, @carlo.bertolli wrote:

This is already a lot of code with parse+sema. I wonder if we should split the patch into two, i.e. 1. parse+sema; 2. code gen? @ABataev ?
It should simplify maintenance of the patch and allow time to extend the OpenMP IR builder.

+1

[tianshilei1992]

I'll separate them later.

[tianshilei1992]

I created a patch D115561 for parser.

clang/include/clang/AST/OpenMPClause.h

const_child_range used_children() const {
     return const_child_range(const_child_iterator(), const_child_iterator());
   }
 
   static bool classof(const OMPClause *T) {
     return T->getClauseKind() == llvm::omp::OMPC_capture;
   }
 };
 
+/// This represents 'compare' clause in the '#pragma omp atomic'
+/// directive.
+///
+/// \code
+/// #pragma omp atomic compare
+/// \endcode
+/// In this example directive '#pragma omp atomic' has 'compare' clause.
+class OMPCompareClause : public OMPClause {
+public:
+  /// Build 'compare' clause.
+  ///
+  /// \param StartLoc Starting location of the clause.
+  /// \param EndLoc Ending location of the clause.
+  OMPCompareClause(SourceLocation StartLoc, SourceLocation EndLoc)
+      : OMPClause(llvm::omp::OMPC_compare, StartLoc, EndLoc) {}
+
+  /// Build an empty clause.
+  OMPCompareClause()
+      : OMPClause(llvm::omp::OMPC_compare, SourceLocation(), SourceLocation()) {
+  }
+
+  child_range children() {
+    return child_range(child_iterator(), child_iterator());
+  }
+
+  const_child_range children() const {
+    return const_child_range(const_child_iterator(), const_child_iterator());
+  }
+
+  child_range used_children() {
+    return child_range(child_iterator(), child_iterator());
+  }
+  const_child_range used_children() const {
+    return const_child_range(const_child_iterator(), const_child_iterator());
+  }
+
+  static bool classof(const OMPClause *T) {
+    return T->getClauseKind() == llvm::omp::OMPC_compare;
+  }
+};
+
 /// This represents 'seq_cst' clause in the '#pragma omp atomic'
 /// directive.
 ///
 /// \code
 /// #pragma omp atomic seq_cst
 /// \endcode
 /// In this example directive '#pragma omp atomic' has 'seq_cst' clause.
 class OMPSeqCstClause : public OMPClause {

clang/include/clang/AST/RecursiveASTVisitor.h

 }
 
 template <typename Derived>
 bool RecursiveASTVisitor<Derived>::VisitOMPCaptureClause(OMPCaptureClause *) {
   return true;
 }
 
 template <typename Derived>
+bool RecursiveASTVisitor<Derived>::VisitOMPCompareClause(OMPCompareClause *) {
+  return true;
+}
+
+template <typename Derived>
 bool RecursiveASTVisitor<Derived>::VisitOMPSeqCstClause(OMPSeqCstClause *) {
   return true;
 }
 
 template <typename Derived>
 bool RecursiveASTVisitor<Derived>::VisitOMPAcqRelClause(OMPAcqRelClause *) {
   return true;
 }

clang/include/clang/AST/StmtOpenMP.h

OMPAtomicDirective(SourceLocation StartLoc, SourceLocation EndLoc)
                                StartLoc, EndLoc) {}
 
   /// Build an empty directive.
   ///
   explicit OMPAtomicDirective()
       : OMPExecutableDirective(OMPAtomicDirectiveClass, llvm::omp::OMPD_atomic,
                                SourceLocation(), SourceLocation()) {}
 
   enum DataPositionTy : size_t {
     POS_X = 0,
     POS_V,
     POS_E,
     POS_UpdateExpr,
+    POS_D,
+    POS_CondExpr
   };
 
   /// Set 'x' part of the associated expression/statement.
   void setX(Expr *X) { Data->getChildren()[DataPositionTy::POS_X] = X; }
+  /// Set 'v' part of the associated expression/statement.
+  void setV(Expr *V) { Data->getChildren()[DataPositionTy::POS_V] = V; }
+  /// Set 'd' part of the associated expression/statement.
+  void setD(Expr *D) { Data->getChildren()[DataPositionTy::POS_D] = D; }
+  /// Set 'e' or 'expr' part of the associated expression/statement.
+  void setE(Expr *E) { Data->getChildren()[DataPositionTy::POS_E] = E; }
   /// Set helper expression of the form
   /// 'OpaqueValueExpr(x) binop OpaqueValueExpr(expr)' or
   /// 'OpaqueValueExpr(expr) binop OpaqueValueExpr(x)'.
   void setUpdateExpr(Expr *UE) {
     Data->getChildren()[DataPositionTy::POS_UpdateExpr] = UE;
   }
-  /// Set 'v' part of the associated expression/statement.
-  void setV(Expr *V) { Data->getChildren()[DataPositionTy::POS_V] = V; }
-  /// Set 'expr' part of the associated expression/statement.
-  void setExpr(Expr *E) { Data->getChildren()[DataPositionTy::POS_E] = E; }
+  /// Set the conditional expression part of the associated expression/statement
+  /// in atomic compare.
+  void setCondExpr(Expr *CE) {
+    Data->getChildren()[DataPositionTy::POS_CondExpr] = CE;
+  }
 
 public:
   /// Creates directive with a list of \a Clauses and 'x', 'v' and 'expr'
   /// parts of the atomic construct (see Section 2.12.6, atomic Construct, for
   /// detailed description of 'x', 'v' and 'expr').
   ///
   /// \param C AST context.
   /// \param StartLoc Starting location of the directive kind.
   /// \param EndLoc Ending Location of the directive.
   /// \param Clauses List of clauses.
   /// \param AssociatedStmt Statement, associated with the directive.
   /// \param X 'x' part of the associated expression/statement.
   /// \param V 'v' part of the associated expression/statement.
   /// \param E 'expr' part of the associated expression/statement.
   /// \param UE Helper expression of the form
   /// 'OpaqueValueExpr(x) binop OpaqueValueExpr(expr)' or
   /// 'OpaqueValueExpr(expr) binop OpaqueValueExpr(x)'.
   /// \param IsXLHSInRHSPart true if \a UE has the first form and false if the
   /// second.
   /// \param IsPostfixUpdate true if original value of 'x' must be stored in
   /// 'v', not an updated one.
   static OMPAtomicDirective *
   Create(const ASTContext &C, SourceLocation StartLoc, SourceLocation EndLoc,
          ArrayRef<OMPClause *> Clauses, Stmt *AssociatedStmt, Expr *X, Expr *V,
-         Expr *E, Expr *UE, bool IsXLHSInRHSPart, bool IsPostfixUpdate);
+         Expr *E, Expr *D, Expr *Ex, Expr *UE, bool IsXLHSInRHSPart,
+         bool IsPostfixUpdate);
 
   /// Creates an empty directive with the place for \a NumClauses
   /// clauses.
   ///
   /// \param C AST context.
   /// \param NumClauses Number of clauses.
   ///
   static OMPAtomicDirective *CreateEmpty(const ASTContext &C,
                                          unsigned NumClauses, EmptyShell);
 
   /// Get 'x' part of the associated expression/statement.
   Expr *getX() {
     return cast_or_null<Expr>(Data->getChildren()[DataPositionTy::POS_X]);
   }
   const Expr *getX() const {
     return cast_or_null<Expr>(Data->getChildren()[DataPositionTy::POS_X]);
   }
+  /// Get 'd' part of the associated expression/statement.
+  Expr *getD() {
+    return cast_or_null<Expr>(Data->getChildren()[DataPositionTy::POS_D]);
+  }
+  const Expr *getD() const {
+    return cast_or_null<Expr>(Data->getChildren()[DataPositionTy::POS_D]);
+  }
+  /// Get 'e' part of the associated expression/statement.
+  Expr *getE() {
+    return cast_or_null<Expr>(Data->getChildren()[DataPositionTy::POS_E]);
+  }
+  const Expr *getE() const {
+    return cast_or_null<Expr>(Data->getChildren()[DataPositionTy::POS_E]);
+  }
+  /// Get 'v' part of the associated expression/statement.
+  Expr *getV() {
+    return cast_or_null<Expr>(Data->getChildren()[DataPositionTy::POS_V]);
+  }
+  const Expr *getV() const {
+    return cast_or_null<Expr>(Data->getChildren()[DataPositionTy::POS_V]);
+  }
   /// Get helper expression of the form
   /// 'OpaqueValueExpr(x) binop OpaqueValueExpr(expr)' or
   /// 'OpaqueValueExpr(expr) binop OpaqueValueExpr(x)'.
   Expr *getUpdateExpr() {
     return cast_or_null<Expr>(
         Data->getChildren()[DataPositionTy::POS_UpdateExpr]);
   }
   const Expr *getUpdateExpr() const {
     return cast_or_null<Expr>(
         Data->getChildren()[DataPositionTy::POS_UpdateExpr]);
   }
+  /// Get the conditional expression.
+  Expr *getCondExpr() {
+    return cast_or_null<Expr>(
+        Data->getChildren()[DataPositionTy::POS_CondExpr]);
+  }
+  const Expr *getCondExpr() const {
+    return cast_or_null<Expr>(
+        Data->getChildren()[DataPositionTy::POS_CondExpr]);
+  }
   /// Return true if helper update expression has form
   /// 'OpaqueValueExpr(x) binop OpaqueValueExpr(expr)' and false if it has form
   /// 'OpaqueValueExpr(expr) binop OpaqueValueExpr(x)'.
   bool isXLHSInRHSPart() const { return IsXLHSInRHSPart; }
   /// Return true if 'v' expression must be updated to original value of
   /// 'x', false if 'v' must be updated to the new value of 'x'.
   bool isPostfixUpdate() const { return IsPostfixUpdate; }
-  /// Get 'v' part of the associated expression/statement.
-  Expr *getV() {
-    return cast_or_null<Expr>(Data->getChildren()[DataPositionTy::POS_V]);
-  }
-  const Expr *getV() const {
-    return cast_or_null<Expr>(Data->getChildren()[DataPositionTy::POS_V]);
-  }
-  /// Get 'expr' part of the associated expression/statement.
-  Expr *getExpr() {
-    return cast_or_null<Expr>(Data->getChildren()[DataPositionTy::POS_E]);
-  }
-  const Expr *getExpr() const {
-    return cast_or_null<Expr>(Data->getChildren()[DataPositionTy::POS_E]);
-  }
 
   static bool classof(const Stmt *T) {
     return T->getStmtClass() == OMPAtomicDirectiveClass;
   }
 };
 
 /// This represents '#pragma omp target' directive.
 ///

clang/include/clang/Basic/DiagnosticSemaKinds.td

 def err_omp_atomic_capture_not_expression_statement : Error<
   "the statement for 'atomic capture' must be an expression statement of form 'v = ++x;', 'v = --x;', 'v = x++;', 'v = x--;', 'v = x binop= expr;', 'v = x = x binop expr' or 'v = x = expr binop x',"
   " where x and v are both lvalue expressions with scalar type">;
 def err_omp_atomic_capture_not_compound_statement : Error<
   "the statement for 'atomic capture' must be a compound statement of form '{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;}' or '{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;}'"
   " where x is an lvalue expression with scalar type">;
+def err_omp_atomic_compare : Error<
+  "the statement for 'atomic compare' must be a compound statement of form '{x = expr ordop x ? expr : x;}', '{x = x ordop expr ? expr : x;}', 'x = x == e ? d : x;',"
+  " '{if (expr ordop x) { x = expr; }}', '{if (x ordop expr) { x = expr; }}', or '{if (x == e) { x = d; }}' where x, v are lvalue expressions with scalar type, e, d are expression with scalar type.">;
 def note_omp_atomic_capture: Note<
   "%select{expected assignment expression|expected compound statement|expected exactly two expression statements|expected in right hand side of the first expression}0">;
+def note_omp_atomic_compare: Note<
+  "%select{expected compound statement of size one|expected conditional update statement|expected assignment expression|expected conditional expression|expected binary operator|expected <, > or == for ordop|expected 'x' at false expression|expected 'x' in conditional expression|expected 'x ordop e'|expected lvalue scalar expression|expected scalar expression|only integers supported}0">;
 def err_omp_atomic_several_clauses : Error<
-  "directive '#pragma omp atomic' cannot contain more than one 'read', 'write', 'update' or 'capture' clause">;
+  "directive '#pragma omp atomic' cannot contain more than one 'read', 'write', 'update', 'capture', or 'compare' clause">;
 def err_omp_several_mem_order_clauses : Error<
   "directive '#pragma omp %0' cannot contain more than one %select{'seq_cst', 'relaxed', |}1'acq_rel', 'acquire' or 'release' clause">;
 def err_omp_atomic_incompatible_mem_order_clause : Error<
   "directive '#pragma omp atomic%select{ %0|}1' cannot be used with '%2' clause">;
 def note_omp_previous_mem_order_clause : Note<
   "'%0' clause used here">;
 def err_omp_target_contains_not_only_teams : Error<
   "target construct with nested teams region contains statements outside of the teams construct">;

clang/include/clang/Sema/Sema.h

public:
   OMPClause *ActOnOpenMPWriteClause(SourceLocation StartLoc,
                                     SourceLocation EndLoc);
   /// Called on well-formed 'update' clause.
   OMPClause *ActOnOpenMPUpdateClause(SourceLocation StartLoc,
                                      SourceLocation EndLoc);
   /// Called on well-formed 'capture' clause.
   OMPClause *ActOnOpenMPCaptureClause(SourceLocation StartLoc,
                                       SourceLocation EndLoc);
+  /// Called on well-formed 'compare' clause.
+  OMPClause *ActOnOpenMPCompareClause(SourceLocation StartLoc,
+                                      SourceLocation EndLoc);
   /// Called on well-formed 'seq_cst' clause.
   OMPClause *ActOnOpenMPSeqCstClause(SourceLocation StartLoc,
                                      SourceLocation EndLoc);
   /// Called on well-formed 'acq_rel' clause.
   OMPClause *ActOnOpenMPAcqRelClause(SourceLocation StartLoc,
                                      SourceLocation EndLoc);
   /// Called on well-formed 'acquire' clause.
   OMPClause *ActOnOpenMPAcquireClause(SourceLocation StartLoc,

clang/lib/AST/OpenMPClause.cpp

const OMPClauseWithPreInit *OMPClauseWithPreInit::get(const OMPClause *C) {
   case OMPC_mergeable:
   case OMPC_threadprivate:
   case OMPC_flush:
   case OMPC_depobj:
   case OMPC_read:
   case OMPC_write:
   case OMPC_update:
   case OMPC_capture:
+  case OMPC_compare:
   case OMPC_seq_cst:
   case OMPC_acq_rel:
   case OMPC_acquire:
   case OMPC_release:
   case OMPC_relaxed:
   case OMPC_depend:
   case OMPC_threads:
   case OMPC_simd:
const OMPClauseWithPostUpdate *OMPClauseWithPostUpdate::get(const OMPClause *C) {
   case OMPC_mergeable:
   case OMPC_threadprivate:
   case OMPC_flush:
   case OMPC_depobj:
   case OMPC_read:
   case OMPC_write:
   case OMPC_update:
   case OMPC_capture:
+  case OMPC_compare:
   case OMPC_seq_cst:
   case OMPC_acq_rel:
   case OMPC_acquire:
   case OMPC_release:
   case OMPC_relaxed:
   case OMPC_depend:
   case OMPC_device:
   case OMPC_threads:
if (Node->isExtended()) {
     OS << ")";
   }
 }
 
 void OMPClausePrinter::VisitOMPCaptureClause(OMPCaptureClause *) {
   OS << "capture";
 }
 
+void OMPClausePrinter::VisitOMPCompareClause(OMPCompareClause *) {
+  OS << "compare";
+}
+
 void OMPClausePrinter::VisitOMPSeqCstClause(OMPSeqCstClause *) {
   OS << "seq_cst";
 }
 
 void OMPClausePrinter::VisitOMPAcqRelClause(OMPAcqRelClause *) {
   OS << "acq_rel";
 }
 

clang/lib/AST/StmtOpenMP.cpp

 OMPOrderedDirective *OMPOrderedDirective::CreateEmpty(const ASTContext &C,
                                                       unsigned NumClauses,
                                                       bool IsStandalone,
                                                       EmptyShell) {
   return createEmptyDirective<OMPOrderedDirective>(C, NumClauses,
                                                    !IsStandalone);
 }
 
-OMPAtomicDirective *OMPAtomicDirective::Create(
-    const ASTContext &C, SourceLocation StartLoc, SourceLocation EndLoc,
-    ArrayRef<OMPClause *> Clauses, Stmt *AssociatedStmt, Expr *X, Expr *V,
-    Expr *E, Expr *UE, bool IsXLHSInRHSPart, bool IsPostfixUpdate) {
+OMPAtomicDirective *
+OMPAtomicDirective::Create(const ASTContext &C, SourceLocation StartLoc,
+                           SourceLocation EndLoc, ArrayRef<OMPClause *> Clauses,
+                           Stmt *AssociatedStmt, Expr *X, Expr *V, Expr *E,
+                           Expr *D, Expr *UE, Expr *CE, bool IsXLHSInRHSPart,
+                           bool IsPostfixUpdate) {
   auto *Dir = createDirective<OMPAtomicDirective>(
-      C, Clauses, AssociatedStmt, /*NumChildren=*/4, StartLoc, EndLoc);
+      C, Clauses, AssociatedStmt, /*NumChildren=*/6, StartLoc, EndLoc);
   Dir->setX(X);
   Dir->setV(V);
-  Dir->setExpr(E);
+  Dir->setE(E);
+  Dir->setD(D);
   Dir->setUpdateExpr(UE);
+  Dir->setCondExpr(CE);
   Dir->IsXLHSInRHSPart = IsXLHSInRHSPart;
   Dir->IsPostfixUpdate = IsPostfixUpdate;
   return Dir;
 }
 
 OMPAtomicDirective *OMPAtomicDirective::CreateEmpty(const ASTContext &C,
                                                     unsigned NumClauses,
                                                     EmptyShell) {

clang/lib/AST/StmtProfile.cpp

 void OMPClauseProfiler::VisitOMPReadClause(const OMPReadClause *) {}
 
 void OMPClauseProfiler::VisitOMPWriteClause(const OMPWriteClause *) {}
 
 void OMPClauseProfiler::VisitOMPUpdateClause(const OMPUpdateClause *) {}
 
 void OMPClauseProfiler::VisitOMPCaptureClause(const OMPCaptureClause *) {}
 
+void OMPClauseProfiler::VisitOMPCompareClause(const OMPCompareClause *) {}
+
 void OMPClauseProfiler::VisitOMPSeqCstClause(const OMPSeqCstClause *) {}
 
 void OMPClauseProfiler::VisitOMPAcqRelClause(const OMPAcqRelClause *) {}
 
 void OMPClauseProfiler::VisitOMPAcquireClause(const OMPAcquireClause *) {}
 
 void OMPClauseProfiler::VisitOMPReleaseClause(const OMPReleaseClause *) {}
 

clang/lib/Basic/OpenMPKinds.cpp

#include "clang/Basic/OpenMPKinds.def"
   case OMPC_nowait:
   case OMPC_untied:
   case OMPC_mergeable:
   case OMPC_flush:
   case OMPC_depobj:
   case OMPC_read:
   case OMPC_write:
   case OMPC_capture:
+  case OMPC_compare:
   case OMPC_seq_cst:
   case OMPC_acq_rel:
   case OMPC_acquire:
   case OMPC_release:
   case OMPC_relaxed:
   case OMPC_threads:
   case OMPC_simd:
   case OMPC_num_teams:
#include "clang/Basic/OpenMPKinds.def"
   case OMPC_nowait:
   case OMPC_untied:
   case OMPC_mergeable:
   case OMPC_flush:
   case OMPC_depobj:
   case OMPC_read:
   case OMPC_write:
   case OMPC_capture:
+  case OMPC_compare:
   case OMPC_seq_cst:
   case OMPC_acq_rel:
   case OMPC_acquire:
   case OMPC_release:
   case OMPC_relaxed:
   case OMPC_threads:
   case OMPC_simd:
   case OMPC_num_teams:

clang/lib/CodeGen/CGStmtOpenMP.cpp

UpdateVal = CGF.Builder.CreateIntCast(
         IC, X.getAddress(CGF).getElementType(),
         X.getType()->hasSignedIntegerRepresentation());
   }
   llvm::Value *Res =
       CGF.Builder.CreateAtomicRMW(RMWOp, X.getPointer(CGF), UpdateVal, AO);
   return std::make_pair(true, RValue::get(Res));
 }
 
+static std::pair<bool, RValue> emitOMPAtomicCmpXchg(CodeGenFunction &CGF,
+                                                    LValue X, RValue E,
+                                                    RValue D,
+                                                    llvm::AtomicOrdering AO) {
+  llvm::Value *Res = CGF.Builder.CreateAtomicCmpXchg(
+      X.getPointer(CGF), E.getScalarVal(), D.getScalarVal(), AO, AO);
+  llvm::Value *Old = CGF.Builder.CreateExtractValue(Res, 0);
+  return std::make_pair(true, RValue::get(Old));
+}
+
 std::pair<bool, RValue> CodeGenFunction::EmitOMPAtomicSimpleUpdateExpr(
     LValue X, RValue E, BinaryOperatorKind BO, bool IsXLHSInRHSPart,
     llvm::AtomicOrdering AO, SourceLocation Loc,
     const llvm::function_ref<RValue(RValue)> CommonGen) {
   // Update expressions are allowed to have the following forms:
   // x binop= expr; -> xrval + expr;
   // x++, ++x -> xrval + 1;
   // x--, --x -> xrval - 1;
case llvm::AtomicOrdering::Monotonic:
       break;
     case llvm::AtomicOrdering::NotAtomic:
     case llvm::AtomicOrdering::Unordered:
       llvm_unreachable("Unexpected ordering.");
     }
   }
 }
 
+static void emitOMPAtomicCompareExpr(CodeGenFunction &CGF,
+                                     llvm::AtomicOrdering AO, const Expr *X,
+                                     const Expr *E, const Expr *D,
+                                     const Expr *CE, bool IsXLHSInRHSPart,
+                                     SourceLocation Loc) {
+  assert(X->isLValue() && "X of 'omp atomic compare' is not lvalue");
+  assert(isa<BinaryOperator>(CE->IgnoreImpCasts()) &&
+         "Cond expr in 'atomic compare' must be a binary operator.");
+
+  auto *CondBO = cast<BinaryOperator>(CE->IgnoreImpCasts());
+
+  LValue XLValue = CGF.EmitLValue(X->IgnoreParenImpCasts());
+  RValue ERValue = CGF.EmitAnyExpr(E->IgnoreParenImpCasts());
+
+  std::pair<bool, RValue> Res;
+
+  if (CondBO->getOpcode() == BO_EQ) {
+    RValue DRValue = CGF.EmitAnyExpr(D->IgnoreParenImpCasts());
+    Res = emitOMPAtomicCmpXchg(CGF, XLValue, ERValue, DRValue, AO);
+  } else
+    Res = emitOMPAtomicRMW(CGF, XLValue, ERValue, CondBO->getOpcode(), AO,
+                           /* IsXLHSInRHSPart */ false);
+  // Cannot emit atomic operation.
+  // TODO: Do we really want to emit a non-atomic operation here?
+  if (!Res.first) {
+  }
+}
+
 static void emitOMPAtomicExpr(CodeGenFunction &CGF, OpenMPClauseKind Kind,
                               llvm::AtomicOrdering AO, bool IsPostfixUpdate,
                               const Expr *X, const Expr *V, const Expr *E,
-                              const Expr *UE, bool IsXLHSInRHSPart,
-                              SourceLocation Loc) {
+                              const Expr *D, const Expr *UE, const Expr *CE,
+                              bool IsXLHSInRHSPart, SourceLocation Loc) {
   switch (Kind) {
   case OMPC_read:
     emitOMPAtomicReadExpr(CGF, AO, X, V, Loc);
     break;
   case OMPC_write:
     emitOMPAtomicWriteExpr(CGF, AO, X, E, Loc);
     break;
   case OMPC_unknown:
   case OMPC_update:
     emitOMPAtomicUpdateExpr(CGF, AO, X, E, UE, IsXLHSInRHSPart, Loc);
     break;
   case OMPC_capture:
     emitOMPAtomicCaptureExpr(CGF, AO, IsPostfixUpdate, V, X, E, UE,
                              IsXLHSInRHSPart, Loc);
     break;
+  case OMPC_compare:
+    emitOMPAtomicCompareExpr(CGF, AO, X, E, D, CE, IsXLHSInRHSPart, Loc);
+    break;
   case OMPC_if:
   case OMPC_final:
   case OMPC_num_threads:
   case OMPC_private:
   case OMPC_firstprivate:
   case OMPC_lastprivate:
   case OMPC_reduction:
   case OMPC_task_reduction:
if (DefaultOrder == llvm::AtomicOrdering::Monotonic ||
         AO = llvm::AtomicOrdering::Acquire;
       }
     }
   }
 
   LexicalScope Scope(*this, S.getSourceRange());
   EmitStopPoint(S.getAssociatedStmt());
   emitOMPAtomicExpr(*this, Kind, AO, S.isPostfixUpdate(), S.getX(), S.getV(),
-                    S.getExpr(), S.getUpdateExpr(), S.isXLHSInRHSPart(),
-                    S.getBeginLoc());
+                    S.getE(), S.getD(), S.getUpdateExpr(), S.getCondExpr(),
+                    S.isXLHSInRHSPart(), S.getBeginLoc());
 }
 
 static void emitCommonOMPTargetDirective(CodeGenFunction &CGF,
                                          const OMPExecutableDirective &S,
                                          const RegionCodeGenTy &CodeGen) {
   assert(isOpenMPTargetExecutionDirective(S.getDirectiveKind()));
   CodeGenModule &CGM = CGF.CGM;
 

clang/lib/Parse/ParseOpenMP.cpp

case OMPC_if:
     Clause = ParseOpenMPSingleExprWithArgClause(DKind, CKind, WrongDirective);
     break;
   case OMPC_nowait:
   case OMPC_untied:
   case OMPC_mergeable:
   case OMPC_read:
   case OMPC_write:
   case OMPC_capture:
+  case OMPC_compare:
   case OMPC_seq_cst:
   case OMPC_acq_rel:
   case OMPC_acquire:
   case OMPC_release:
   case OMPC_relaxed:
   case OMPC_threads:
   case OMPC_simd:
   case OMPC_nogroup:

clang/lib/Sema/SemaOpenMP.cpp

for (OMPClause *C : Clauses) {
       case OMPC_nowait:
       case OMPC_untied:
       case OMPC_mergeable:
       case OMPC_allocate:
       case OMPC_read:
       case OMPC_write:
       case OMPC_update:
       case OMPC_capture:
+      case OMPC_compare:
       case OMPC_seq_cst:
       case OMPC_acq_rel:
       case OMPC_acquire:
       case OMPC_release:
       case OMPC_relaxed:
       case OMPC_depend:
       case OMPC_threads:
       case OMPC_simd:
public:
   /// true if the source expression is a postfix unary operation, false
   /// if it is a prefix unary operation.
   bool isPostfixUpdate() const { return IsPostfixUpdate; }
 
 private:
   bool checkBinaryOperation(BinaryOperator *AtomicBinOp, unsigned DiagId = 0,
                             unsigned NoteId = 0);
 };
+
+/// Get the node id of the fixed point of corresponding statement
+llvm::FoldingSetNodeID getNodeId(ASTContext &Context, const Expr *S) {
+  llvm::FoldingSetNodeID Id;
+  S->IgnoreParenImpCasts()->Profile(Id, Context, true);
+  return Id;
+}
+
+/// Check if two statemenmts are the same fixed point statements
+bool checkIfTwoExprsAreSame(ASTContext &Context, const Expr *LHS,
+                            const Expr *RHS) {
+  return getNodeId(Context, LHS) == getNodeId(Context, RHS);
+}
+
+// OpenMP 5.1 [2.19.7 atomic Construct]
+// cond-update-stmt, a conditional update statement that has one of the
+// following forms:
+// if (expr ordop x) { x = expr; }
+// if (x ordop expr) { x = expr; }
+// if (x == e) { x = d; }
+bool checkIfCondUpdateStmt(ASTContext &Context, IfStmt *IS, Expr *&X, Expr *&E,
+                           Expr *&D, Expr *&C) {
+  auto *Cond = dyn_cast<BinaryOperator>(IS->getCond());
+  auto *Then = dyn_cast<BinaryOperator>(IS->getThen());
+  // Then can be compound statement
+  if (!Then && isa<CompoundStmt>(IS->getThen())) {
+    auto *CS = cast<CompoundStmt>(IS->getThen());
+    if (CS->size() != 1)
+      return false;
+    Then = dyn_cast<BinaryOperator>(*CS->child_begin());
+  }
+
+  if (!Cond || !Then)
+    return false;
+
+  // The if-stmt cannot have else statement
+  if (IS->getElse())
+    return false;
+
+  // Only ==, <, > are supported
+  if (Cond->getOpcode() != BO_EQ && Cond->getOpcode() != BO_LT &&
+      Cond->getOpcode() != BO_GT)
+    return false;
+
+  X = Then->getLHS();
+  C = Cond;
+
+  llvm::FoldingSetNodeID XId = getNodeId(Context, X);
+  llvm::FoldingSetNodeID LHSId = getNodeId(Context, Cond->getLHS());
+  llvm::FoldingSetNodeID RHSId = getNodeId(Context, Cond->getRHS());
+
+  if (Cond->getOpcode() == BO_EQ) {
+    // if (x == e) { x = d; }
+    D = Then->getRHS();
+    if (LHSId == XId)
+      E = Cond->getRHS();
+    else if (RHSId == XId)
+      E = Cond->getLHS();
+    else
+      return false;
+  } else {
+    // if (expr ordop x) { x = expr; }
+    // if (x ordop expr) { x = expr; }
+    E = Then->getRHS();
+
+    llvm::FoldingSetNodeID EId = getNodeId(Context, E);
+    if (LHSId == EId) {
+      if (RHSId != XId)
+        return false;
+    } else if (RHSId == EId) {
+      if (LHSId != XId)
+        return false;
+    } else
+      return false;
+  }
+
+  return true;
+}
+
+/// Helper class for checking expression in 'omp atomic compare' construct.
+class OpenMPAtomicCompareChecker {
+protected:
+  Sema &SemaRef;
+  ASTContext &Context;
+  /// 'x' lvalue part of the source atomic expression.
+  Expr *X = nullptr;
+  /// Depending on different cases, it can be:
+  /// 'e' rvalue part of the source atomic expression.
+  /// 'expr' rvalue part of the source atomic expression.
+  Expr *E = nullptr;
+  /// 'd' rvalue part of the source atomic expression.
+  Expr *D = nullptr;
+  /// 'cond' rvalue part of the source atomic expression, which is in one of the
+  /// following forms:
+  /// expr ordop x
+  /// x ordop expr
+  /// x == e
+  /// Note that when 'x' is on RHS, a corresponding LHS version is generated for
+  /// easier codegen.
+  Expr *C = nullptr;
+
+public:
+  /// All kinds of errors that can occur in `atomic compare`
+  enum class ErrorTy {
+    /// Not only one statement in compound statement
+    NotOneSubstatement = 0,
+    /// Not a cond-update statement (ref. spec 5.1)
+    NotCondUpdateStatement,
+    /// Not an assignment statement
+    NotAssignmentStatement,
+    /// Not a conditional statement
+    NotCondStatement,
+    /// Not a binary operator
+    NotBinaryOperator,
+    /// Not required ordop (ordop must be <, >, or ==)
+    NotRequiredOrderOp,
+    /// 'x' is not at false expression
+    NotAtFalseExpr,
+    /// 'x' is not in conditional expression
+    NotInCondExpr,
+    /// Not required expression 'x ordop e' or 'e ordop x'
+    NotAOrdopEExpr,
+    /// Not a lvalue scalar type expression
+    NotLValueScalarExpr,
+    /// Not a scalar type expression
+    NotScalarExpr,
+    /// Not an integral expression
+    NotIntegerExpr,
+    /// No error
+    NoError
+  };
+
+  /// Error descriptor type which will be returned to Sema
+  struct ErrorDescTy {
+    ErrorTy ErrorNo = ErrorTy::NoError;
+    SourceLocation ErrorLoc, NoteLoc;
+    SourceRange ErrorRange, NoteRange;
+  };
+
+  OpenMPAtomicCompareChecker(Sema &S)
+      : SemaRef(S), Context(S.getASTContext()) {}
+
+  /// Note: If `checkStatement` returns false, it is undefined behavior to call
+  /// the following getters.
+  /// Return the 'x' lvalue part of the source atomic expression.
+  Expr *getX() const { return X; }
+  /// Return the 'e' or 'expr' rvalue part of the source atomic expression.
+  Expr *getE() const { return E; }
+  /// Return the 'd' rvalue part of the source atomic expression.
+  Expr *getD() const { return D; }
+  /// Return the 'cond' rvalue part of the source atomic expression.
+  Expr *getCondExpr() const { return C; }
+
+  /// Check given statement if it conforms with requiremenet of `atomic
+  /// compare`, and set `x`, `e`, `d`, `expr` and conditional expr accordingly.
+  /// if (expr ordop x) { x = expr; }
+  /// if (x ordop expr) { x = expr; }
+  /// if (x == e) { x = d; }
+  /// x = expr ordop x ? expr : x;
+  /// x = expr ordop x ? expr : x;
+  /// x = x == e ? d : x;
+  bool checkStatement(Stmt *S);
+
+  /// Return the error descriptor that will guide the error message emission.
+  /// Note: it is undefined behavior to call it if `checkStatement` returns
+  /// true.
+  ErrorDescTy getErrorDesc() const { return ErrorDesc; }
+
+protected:
+  /// Error descriptor
+  ErrorDescTy ErrorDesc;
+  /// Check if all results conform with spec in terms of lvalue/rvalue and
+  /// scalar type.
+  bool checkResult();
+  /// Update {expr ordop x} to {x ordop expr} for easier codegen if needed.
+  void updateCondExpr();
+};
 } // namespace
 
+void OpenMPAtomicCompareChecker::updateCondExpr() {
+  // We don't apply any correctness check in this function.
+  auto *Cond = cast<BinaryOperator>(C);
+
+  if (Cond->getOpcode() == BO_EQ)
+    return;
+
+  if (checkIfTwoExprsAreSame(Context, X, Cond->getLHS()))
+    return;
+
+  BinaryOperatorKind Op = Cond->getOpcode() == BO_LT ? BO_GT : BO_LT;
+  ExprResult Result =
+      SemaRef.CreateBuiltinBinOp(Cond->getOperatorLoc(), Op, X, E);
+  assert(!Result.isInvalid());
+
+  C = Result.get();
+}
+
+bool OpenMPAtomicCompareChecker::checkStatement(Stmt *S) {
+  auto *CS = dyn_cast<CompoundStmt>(S);
+  if (CS) {
+    if (CS->size() != 1) {
+      ErrorDesc.ErrorNo = ErrorTy::NotOneSubstatement;
+      ErrorDesc.NoteRange = ErrorDesc.ErrorRange = CS->getSourceRange();
+      ErrorDesc.NoteLoc = ErrorDesc.ErrorLoc = CS->getBeginLoc();
+      return false;
+    }
+    S = CS->body_front();
+  }
+
+  // Check if the statement is in one of the following forms (cond-update-stmt):
+  // if (expr ordop x) { x = expr; }
+  // if (x ordop expr) { x = expr; }
+  // if (x == e) { x = d; }
+  if (auto *IS = dyn_cast<IfStmt>(S)) {
+    if (!checkIfCondUpdateStmt(Context, IS, X, E, D, C)) {
+      ErrorDesc.ErrorNo = ErrorTy::NotCondUpdateStatement;
+      ErrorDesc.NoteRange = ErrorDesc.ErrorRange = IS->getSourceRange();
+      ErrorDesc.NoteLoc = ErrorDesc.ErrorLoc = IS->getBeginLoc();
+      return false;
+    }
+
+    updateCondExpr();
+
+    return checkResult();
+  }
+
+  // Check if the statement is in one of the following forms (cond-expr-stmt):
+  // x = expr ordop x ? expr : x;
+  // x = x ordop expr ? expr : x;
+  // x = x == e ? d : x;
+  auto *BO = dyn_cast<BinaryOperator>(S);
+  if (!BO || BO->getOpcode() != BO_Assign) {
+    ErrorDesc.ErrorNo = ErrorTy::NotAssignmentStatement;
+    ErrorDesc.NoteRange = ErrorDesc.ErrorRange = S->getSourceRange();
+    ErrorDesc.NoteLoc = ErrorDesc.ErrorLoc = S->getBeginLoc();
+    return false;
+  }
+
+  X = BO->getLHS();
+
+  auto *CO = dyn_cast<ConditionalOperator>(BO->getRHS()->IgnoreParenImpCasts());
+  if (!CO) {
+    ErrorDesc.ErrorNo = ErrorTy::NotCondStatement;
+    ErrorDesc.NoteRange = ErrorDesc.ErrorRange = BO->getRHS()->getSourceRange();
+    ErrorDesc.NoteLoc = ErrorDesc.ErrorLoc = BO->getRHS()->getBeginLoc();
+    return false;
+  }
+
+  if (!checkIfTwoExprsAreSame(Context, X, CO->getFalseExpr())) {
+    ErrorDesc.ErrorNo = ErrorTy::NotAtFalseExpr;
+    ErrorDesc.NoteRange = ErrorDesc.ErrorRange =
+        CO->getFalseExpr()->getSourceRange();
+    ErrorDesc.NoteLoc = ErrorDesc.ErrorLoc = CO->getFalseExpr()->getExprLoc();
+    return false;
+  }
+
+  auto *Cond = dyn_cast<BinaryOperator>(CO->getCond());
+  if (!Cond) {
+    ErrorDesc.ErrorNo = ErrorTy::NotBinaryOperator;
+    ErrorDesc.NoteRange = ErrorDesc.ErrorRange =
+        CO->getCond()->getSourceRange();
+    ErrorDesc.NoteLoc = ErrorDesc.ErrorLoc = CO->getCond()->getBeginLoc();
+    return false;
+  }
+
+  // Only ==, <, > are supported
+  if (Cond->getOpcode() != BO_LT && Cond->getOpcode() != BO_GT &&
+      Cond->getOpcode() != BO_EQ) {
+    ErrorDesc.ErrorNo = ErrorTy::NotRequiredOrderOp;
+    ErrorDesc.NoteRange = ErrorDesc.ErrorRange = Cond->getSourceRange();
+    ErrorDesc.NoteLoc = ErrorDesc.ErrorLoc = Cond->getOperatorLoc();
+    return false;
+  }
+
+  if (Cond->getOpcode() == BO_EQ) {
+    if (checkIfTwoExprsAreSame(Context, X, Cond->getLHS()))
+      E = Cond->getRHS();
+    else if (checkIfTwoExprsAreSame(Context, X, Cond->getRHS()))
+      E = Cond->getLHS();
+    else {
+      ErrorDesc.ErrorNo = ErrorTy::NotInCondExpr;
+      ErrorDesc.NoteRange = ErrorDesc.ErrorRange = Cond->getSourceRange();
+      ErrorDesc.NoteLoc = ErrorDesc.ErrorLoc = Cond->getExprLoc();
+      return false;
+    }
+    D = CO->getTrueExpr();
+  } else {
+    E = CO->getTrueExpr();
+    if (!((checkIfTwoExprsAreSame(Context, X, Cond->getLHS()) &&
+           checkIfTwoExprsAreSame(Context, E, Cond->getRHS())) ||
+          (checkIfTwoExprsAreSame(Context, X, Cond->getRHS()) &&
+           checkIfTwoExprsAreSame(Context, E, Cond->getLHS())))) {
+      ErrorDesc.ErrorNo = ErrorTy::NotAOrdopEExpr;
+      ErrorDesc.ErrorRange = Cond->getSourceRange();
+      ErrorDesc.NoteRange = X->getSourceRange();
+      ErrorDesc.ErrorLoc = Cond->getExprLoc();
+      ErrorDesc.NoteLoc = X->getExprLoc();
+      return false;
+    }
+  }
+
+  C = Cond;
+
+  updateCondExpr();
+
+  return checkResult();
+}
+
+bool OpenMPAtomicCompareChecker::checkResult() {
+  // 'x' and 'e' cannot be nullptr
+  assert(X && E && "X and E cannot be nullptr");
+
+  // 'x' and 'e' have to be scalar type. 'x' has to be lvalue, but no such this
+  // requirement for 'e' and 'd'.
+  if (!X->isLValue() || !X->getType()->isScalarType()) {
+    ErrorDesc.ErrorNo = ErrorTy::NotLValueScalarExpr;
+    ErrorDesc.NoteLoc = ErrorDesc.ErrorLoc = X->getExprLoc();
+    ErrorDesc.ErrorRange = ErrorDesc.NoteRange = X->getSourceRange();
+    return false;
+  }
+
+  if (!E->getType()->isScalarType()) {
+    ErrorDesc.ErrorNo = ErrorTy::NotScalarExpr;
+    ErrorDesc.NoteLoc = ErrorDesc.ErrorLoc = E->getExprLoc();
+    ErrorDesc.ErrorRange = ErrorDesc.NoteRange = E->getSourceRange();
+    return false;
+  }
+
+  if (D && !D->getType()->isScalarType()) {
+    ErrorDesc.ErrorNo = ErrorTy::NotScalarExpr;
+    ErrorDesc.NoteLoc = ErrorDesc.ErrorLoc = D->getExprLoc();
+    ErrorDesc.ErrorRange = ErrorDesc.NoteRange = D->getSourceRange();
+    return false;
+  }
+
+  // Currently we only support integer type in CodeGen. We check the
+  // compatibility here. The following code can be removed if we figure out how
+  // to support floating-point types.
+  auto &&CheckIfInteger = [&](Expr *E) {
+    if (!E->getType()->isIntegerType()) {
+      ErrorDesc.ErrorNo = ErrorTy::NotIntegerExpr;
+      ErrorDesc.NoteLoc = ErrorDesc.ErrorLoc = E->getExprLoc();
+      ErrorDesc.ErrorRange = ErrorDesc.NoteRange = E->getSourceRange();
+      return false;
+    }
+
+    return true;
+  };
+
+  if (!CheckIfInteger(X))
+    return false;
+
+  if (!CheckIfInteger(E))
+    return false;
+
+  if (D && !CheckIfInteger(D))
+    return false;
+
+  return true;
+}
+
 bool OpenMPAtomicUpdateChecker::checkBinaryOperation(
     BinaryOperator *AtomicBinOp, unsigned DiagId, unsigned NoteId) {
   ExprAnalysisErrorCode ErrorFound = NoError;
   SourceLocation ErrorLoc, NoteLoc;
   SourceRange ErrorRange, NoteRange;
   // Allowed constructs are:
   //  x = x binop expr;
   //  x = expr binop x;
StmtResult Sema::ActOnOpenMPAtomicDirective(ArrayRef<OMPClause *> Clauses,
   // top and a single exit at the bottom.
   // The point of exit cannot be a branch out of the structured block.
   // longjmp() and throw() must not violate the entry/exit criteria.
   OpenMPClauseKind AtomicKind = OMPC_unknown;
   SourceLocation AtomicKindLoc;
   OpenMPClauseKind MemOrderKind = OMPC_unknown;
   SourceLocation MemOrderLoc;
   for (const OMPClause *C : Clauses) {
-    if (C->getClauseKind() == OMPC_read || C->getClauseKind() == OMPC_write ||
-        C->getClauseKind() == OMPC_update ||
-        C->getClauseKind() == OMPC_capture) {
-      if (AtomicKind != OMPC_unknown) {
+    switch (C->getClauseKind()) {
+    default:
+      llvm_unreachable("unknow clause for atomic directive");
+    case OMPC_read:
+    case OMPC_write:
+    case OMPC_update:
+    case OMPC_capture:
+    case OMPC_compare:
+      if (AtomicKind == OMPC_unknown) {
+        AtomicKind = C->getClauseKind();
+        AtomicKindLoc = C->getBeginLoc();
+        break;
+      } else {
         Diag(C->getBeginLoc(), diag::err_omp_atomic_several_clauses)
             << SourceRange(C->getBeginLoc(), C->getEndLoc());
         Diag(AtomicKindLoc, diag::note_omp_previous_mem_order_clause)
             << getOpenMPClauseName(AtomicKind);
-      } else {
-        AtomicKind = C->getClauseKind();
-        AtomicKindLoc = C->getBeginLoc();
-      }
+        return StmtError();
       }
-    if (C->getClauseKind() == OMPC_seq_cst ||
-        C->getClauseKind() == OMPC_acq_rel ||
-        C->getClauseKind() == OMPC_acquire ||
-        C->getClauseKind() == OMPC_release ||
-        C->getClauseKind() == OMPC_relaxed) {
+    case OMPC_seq_cst:
+    case OMPC_acq_rel:
+    case OMPC_acquire:
+    case OMPC_release:
+    case OMPC_relaxed:
       if (MemOrderKind != OMPC_unknown) {
         Diag(C->getBeginLoc(), diag::err_omp_several_mem_order_clauses)
             << getOpenMPDirectiveName(OMPD_atomic) << 0
             << SourceRange(C->getBeginLoc(), C->getEndLoc());
         Diag(MemOrderLoc, diag::note_omp_previous_mem_order_clause)
             << getOpenMPClauseName(MemOrderKind);
       } else {
         MemOrderKind = C->getClauseKind();
         MemOrderLoc = C->getBeginLoc();
       }
+      break;
     }
   }
+
   // OpenMP 5.0, 2.17.7 atomic Construct, Restrictions
   // If atomic-clause is read then memory-order-clause must not be acq_rel or
   // release.
   // If atomic-clause is write then memory-order-clause must not be acq_rel or
   // acquire.
   // If atomic-clause is update or not present then memory-order-clause must not
   // be acq_rel or acquire.
   if ((AtomicKind == OMPC_read &&
if ((AtomicKind == OMPC_read &&
     Diag(MemOrderLoc, diag::note_omp_previous_mem_order_clause)
         << getOpenMPClauseName(MemOrderKind);
   }
 
   Stmt *Body = AStmt;
   if (auto *EWC = dyn_cast<ExprWithCleanups>(Body))
     Body = EWC->getSubExpr();
 
+  // Stands for 'x' in the spec
   Expr *X = nullptr;
+  // Stands for 'v' in the spec
   Expr *V = nullptr;
+  // Stands for 'd' in the spec
+  Expr *D = nullptr;
+  // Stands for 'e' or 'expr' in the spec
   Expr *E = nullptr;
+  // Stands for update-stmt in the spec
   Expr *UE = nullptr;
+  // Stands for conditional statement in the spec
+  Expr *CE = nullptr;
   bool IsXLHSInRHSPart = false;
   bool IsPostfixUpdate = false;
   // OpenMP [2.12.6, atomic Construct]
   // In the next expressions:
   // * x and v (as applicable) are both l-value expressions with scalar type.
   // * During the execution of an atomic region, multiple syntactic
   // occurrences of x must designate the same storage location.
   // * Neither of v and expr (as applicable) may access the storage location
if (AtomicKind == OMPC_read) {
     } else {
       ErrorFound = NotAnExpression;
       NoteLoc = ErrorLoc = Body->getBeginLoc();
       NoteRange = ErrorRange = SourceRange(NoteLoc, NoteLoc);
     }
     if (ErrorFound != NoError) {
       Diag(ErrorLoc, diag::err_omp_atomic_read_not_expression_statement)
           << ErrorRange;
-      Diag(NoteLoc, diag::note_omp_atomic_read_write) << ErrorFound
-                                                      << NoteRange;
+      Diag(NoteLoc, diag::note_omp_atomic_read_write)
+          << ErrorFound << NoteRange;
       return StmtError();
     }
     if (CurContext->isDependentContext())
       V = X = nullptr;
   } else if (AtomicKind == OMPC_write) {
     enum {
       NotAnExpression,
       NotAnAssignmentOp,
if (AtomicKind == OMPC_read) {
     } else {
       ErrorFound = NotAnExpression;
       NoteLoc = ErrorLoc = Body->getBeginLoc();
       NoteRange = ErrorRange = SourceRange(NoteLoc, NoteLoc);
     }
     if (ErrorFound != NoError) {
       Diag(ErrorLoc, diag::err_omp_atomic_write_not_expression_statement)
           << ErrorRange;
-      Diag(NoteLoc, diag::note_omp_atomic_read_write) << ErrorFound
-                                                      << NoteRange;
+      Diag(NoteLoc, diag::note_omp_atomic_read_write)
+          << ErrorFound << NoteRange;
       return StmtError();
     }
     if (CurContext->isDependentContext())
       E = X = nullptr;
   } else if (AtomicKind == OMPC_update || AtomicKind == OMPC_unknown) {
     // If clause is update:
     //  x++;
     //  x--;
     //  ++x;
     //  --x;
     //  x binop= expr;
     //  x = x binop expr;
     //  x = expr binop x;
     OpenMPAtomicUpdateChecker Checker(*this);
     if (Checker.checkStatement(
-            Body, (AtomicKind == OMPC_update)
-                      ? diag::err_omp_atomic_update_not_expression_statement
-                      : diag::err_omp_atomic_not_expression_statement,
+            Body,
+            (AtomicKind == OMPC_update)
             diag::note_omp_atomic_update))
       return StmtError();
     if (!CurContext->isDependentContext()) {
       E = Checker.getExpr();
       X = Checker.getX();
       UE = Checker.getUpdateExpr();
       IsXLHSInRHSPart = Checker.isXLHSInRHSPart();
     }
if (const auto *AtomicBody = dyn_cast<Expr>(Body)) {
           ErrorFound = NotTwoSubstatements;
         }
       } else {
         NoteLoc = ErrorLoc = Body->getBeginLoc();
         NoteRange = ErrorRange =
             SourceRange(Body->getBeginLoc(), Body->getBeginLoc());
         ErrorFound = NotACompoundStatement;
       }
-      if (ErrorFound != NoError) {
-        Diag(ErrorLoc, diag::err_omp_atomic_capture_not_compound_statement)
-            << ErrorRange;
-        Diag(NoteLoc, diag::note_omp_atomic_capture) << ErrorFound << NoteRange;
-        return StmtError();
-      }
-      if (CurContext->isDependentContext())
-        UE = V = E = X = nullptr;
+    }
+  } else if (AtomicKind == OMPC_compare) {
+    // TODO: error handling
+    OpenMPAtomicCompareChecker Checker(*this);
+    if (!Checker.checkStatement(Body)) {
+      OpenMPAtomicCompareChecker::ErrorDescTy Err = Checker.getErrorDesc();
+      Diag(Err.ErrorLoc, diag::err_omp_atomic_compare) << Err.ErrorRange;
+      Diag(Err.NoteLoc, diag::note_omp_atomic_compare)
+          << static_cast<int>(Err.ErrorNo) << Err.NoteRange;
+      return StmtError();
     }
+    X = Checker.getX();
+    E = Checker.getE();
+    D = Checker.getD();
+    CE = Checker.getCondExpr();
+    if (CurContext->isDependentContext())
+      X = E = D = CE = nullptr;
   }
 
   setFunctionHasBranchProtectedScope();
 
   return OMPAtomicDirective::Create(Context, StartLoc, EndLoc, Clauses, AStmt,
-                                    X, V, E, UE, IsXLHSInRHSPart,
+                                    X, V, E, D, UE, CE, IsXLHSInRHSPart,
                                     IsPostfixUpdate);
 }
 
 StmtResult Sema::ActOnOpenMPTargetDirective(ArrayRef<OMPClause *> Clauses,
                                             Stmt *AStmt,
                                             SourceLocation StartLoc,
                                             SourceLocation EndLoc) {
   if (!AStmt)
OMPClause *Sema::ActOnOpenMPSingleExprClause(OpenMPClauseKind Kind, Expr *Expr,
   case OMPC_threadprivate:
   case OMPC_sizes:
   case OMPC_allocate:
   case OMPC_flush:
   case OMPC_read:
   case OMPC_write:
   case OMPC_update:
   case OMPC_capture:
+  case OMPC_compare:
   case OMPC_seq_cst:
   case OMPC_acq_rel:
   case OMPC_acquire:
   case OMPC_release:
   case OMPC_relaxed:
   case OMPC_depend:
   case OMPC_threads:
   case OMPC_simd:
static OpenMPDirectiveKind getOpenMPCaptureRegionForClause(
   case OMPC_threadprivate:
   case OMPC_allocate:
   case OMPC_flush:
   case OMPC_depobj:
   case OMPC_read:
   case OMPC_write:
   case OMPC_update:
   case OMPC_capture:
+  case OMPC_compare:
   case OMPC_seq_cst:
   case OMPC_acq_rel:
   case OMPC_acquire:
   case OMPC_release:
   case OMPC_relaxed:
   case OMPC_depend:
   case OMPC_threads:
   case OMPC_simd:
OMPClause *Sema::ActOnOpenMPSimpleClause(
   case OMPC_mergeable:
   case OMPC_threadprivate:
   case OMPC_allocate:
   case OMPC_flush:
   case OMPC_depobj:
   case OMPC_read:
   case OMPC_write:
   case OMPC_capture:
+  case OMPC_compare:
   case OMPC_seq_cst:
   case OMPC_acq_rel:
   case OMPC_acquire:
   case OMPC_release:
   case OMPC_relaxed:
   case OMPC_depend:
   case OMPC_device:
   case OMPC_threads:
OMPClause *Sema::ActOnOpenMPSingleExprWithArgClause(
   case OMPC_threadprivate:
   case OMPC_allocate:
   case OMPC_flush:
   case OMPC_depobj:
   case OMPC_read:
   case OMPC_write:
   case OMPC_update:
   case OMPC_capture:
+  case OMPC_compare:
   case OMPC_seq_cst:
   case OMPC_acq_rel:
   case OMPC_acquire:
   case OMPC_release:
   case OMPC_relaxed:
   case OMPC_depend:
   case OMPC_threads:
   case OMPC_simd:
case OMPC_write:
     Res = ActOnOpenMPWriteClause(StartLoc, EndLoc);
     break;
   case OMPC_update:
     Res = ActOnOpenMPUpdateClause(StartLoc, EndLoc);
     break;
   case OMPC_capture:
     Res = ActOnOpenMPCaptureClause(StartLoc, EndLoc);
     break;
+  case OMPC_compare:
+    Res = ActOnOpenMPCompareClause(StartLoc, EndLoc);
+    break;
   case OMPC_seq_cst:
     Res = ActOnOpenMPSeqCstClause(StartLoc, EndLoc);
     break;
   case OMPC_acq_rel:
     Res = ActOnOpenMPAcqRelClause(StartLoc, EndLoc);
     break;
   case OMPC_acquire:
     Res = ActOnOpenMPAcquireClause(StartLoc, EndLoc);
OMPClause *Sema::ActOnOpenMPUpdateClause(SourceLocation StartLoc,
   return OMPUpdateClause::Create(Context, StartLoc, EndLoc);
 }
 
 OMPClause *Sema::ActOnOpenMPCaptureClause(SourceLocation StartLoc,
                                           SourceLocation EndLoc) {
   return new (Context) OMPCaptureClause(StartLoc, EndLoc);
 }
 
+OMPClause *Sema::ActOnOpenMPCompareClause(SourceLocation StartLoc,
+                                          SourceLocation EndLoc) {
+  return new (Context) OMPCompareClause(StartLoc, EndLoc);
+}
+
 OMPClause *Sema::ActOnOpenMPSeqCstClause(SourceLocation StartLoc,
                                          SourceLocation EndLoc) {
   return new (Context) OMPSeqCstClause(StartLoc, EndLoc);
 }
 
 OMPClause *Sema::ActOnOpenMPAcqRelClause(SourceLocation StartLoc,
                                          SourceLocation EndLoc) {
   return new (Context) OMPAcqRelClause(StartLoc, EndLoc);
OMPClause *Sema::ActOnOpenMPVarListClause(
   case OMPC_nowait:
   case OMPC_untied:
   case OMPC_mergeable:
   case OMPC_threadprivate:
   case OMPC_read:
   case OMPC_write:
   case OMPC_update:
   case OMPC_capture:
+  case OMPC_compare:
   case OMPC_seq_cst:
   case OMPC_acq_rel:
   case OMPC_acquire:
   case OMPC_release:
   case OMPC_relaxed:
   case OMPC_device:
   case OMPC_threads:
   case OMPC_simd:

clang/lib/Sema/TreeTransform.h

 OMPClause *
 TreeTransform<Derived>::TransformOMPCaptureClause(OMPCaptureClause *C) {
   // No need to rebuild this clause, no template-dependent parameters.
   return C;
 }
 
 template <typename Derived>
 OMPClause *
+TreeTransform<Derived>::TransformOMPCompareClause(OMPCompareClause *C) {
+  // No need to rebuild this clause, no template-dependent parameters.
+  return C;
+}
+
+template <typename Derived>
+OMPClause *
 TreeTransform<Derived>::TransformOMPSeqCstClause(OMPSeqCstClause *C) {
   // No need to rebuild this clause, no template-dependent parameters.
   return C;
 }
 
 template <typename Derived>
 OMPClause *
 TreeTransform<Derived>::TransformOMPAcqRelClause(OMPAcqRelClause *C) {

clang/lib/Serialization/ASTReader.cpp

case llvm::omp::OMPC_write:
     C = new (Context) OMPWriteClause();
     break;
   case llvm::omp::OMPC_update:
     C = OMPUpdateClause::CreateEmpty(Context, Record.readInt());
     break;
   case llvm::omp::OMPC_capture:
     C = new (Context) OMPCaptureClause();
     break;
+  case llvm::omp::OMPC_compare:
+    C = new (Context) OMPCompareClause();
+    break;
   case llvm::omp::OMPC_seq_cst:
     C = new (Context) OMPSeqCstClause();
     break;
   case llvm::omp::OMPC_acq_rel:
     C = new (Context) OMPAcqRelClause();
     break;
   case llvm::omp::OMPC_acquire:
     C = new (Context) OMPAcquireClause();
if (C->isExtended()) {
     C->setLParenLoc(Record.readSourceLocation());
     C->setArgumentLoc(Record.readSourceLocation());
     C->setDependencyKind(Record.readEnum<OpenMPDependClauseKind>());
   }
 }
 
 void OMPClauseReader::VisitOMPCaptureClause(OMPCaptureClause *) {}
 
+void OMPClauseReader::VisitOMPCompareClause(OMPCompareClause *) {}
+
 void OMPClauseReader::VisitOMPSeqCstClause(OMPSeqCstClause *) {}
 
 void OMPClauseReader::VisitOMPAcqRelClause(OMPAcqRelClause *) {}
 
 void OMPClauseReader::VisitOMPAcquireClause(OMPAcquireClause *) {}
 
 void OMPClauseReader::VisitOMPReleaseClause(OMPReleaseClause *) {}
 

clang/lib/Serialization/ASTWriter.cpp

if (C->isExtended()) {
     Record.AddSourceLocation(C->getLParenLoc());
     Record.AddSourceLocation(C->getArgumentLoc());
     Record.writeEnum(C->getDependencyKind());
   }
 }
 
 void OMPClauseWriter::VisitOMPCaptureClause(OMPCaptureClause *) {}
 
+void OMPClauseWriter::VisitOMPCompareClause(OMPCompareClause *) {}
+
 void OMPClauseWriter::VisitOMPSeqCstClause(OMPSeqCstClause *) {}
 
 void OMPClauseWriter::VisitOMPAcqRelClause(OMPAcqRelClause *) {}
 
 void OMPClauseWriter::VisitOMPAcquireClause(OMPAcquireClause *) {}
 
 void OMPClauseWriter::VisitOMPReleaseClause(OMPReleaseClause *) {}
 

clang/test/OpenMP/atomic_compare_codegen.cpp

+
+// RUN: %clang_cc1 -verify -triple x86_64-apple-darwin10 -target-cpu core2 -fopenmp -fopenmp-version=51 -x c -emit-llvm %s -o - | FileCheck %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=51 -x c -triple x86_64-apple-darwin10 -target-cpu core2 -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=51 -x c -triple x86_64-apple-darwin10 -target-cpu core2 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s
+
+// RUN: %clang_cc1 -verify -triple x86_64-apple-darwin10 -target-cpu core2 -fopenmp-simd -fopenmp-version=51 -x c -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=51 -x c -triple x86_64-apple-darwin10 -target-cpu core2 -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=51 -x c -triple x86_64-apple-darwin10 -target-cpu core2 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
+// expected-no-diagnostics
+
+// Integral types
+char cx, ce, cd;
+unsigned char ucx, uce, ucd;
+short sx, se, sd;
+unsigned short usx, use, usd;
+int ix, ie, id;
+unsigned int uix, uie, uid;
+long lx, le, ld;
+unsigned long ulx, ule, uld;
+long long llx, lle, lld;
+unsigned long long ullx, ulle, ulld;
+
+void foo() {
+// char
+// {
+
+#pragma omp atomic compare
+  cx = ce > cx ? ce : cx;
+
+#pragma omp atomic compare
+  cx = ce < cx ? ce : cx;
+
+#pragma omp atomic compare
+  cx = cx > ce ? ce : cx;
+
+#pragma omp atomic compare
+  cx = cx < ce ? ce : cx;
+
+#pragma omp atomic compare
+  cx = cx == ce ? cd : cx;
+
+#pragma omp atomic compare
+  if (ce > cx) { cx = ce; }
+
+#pragma omp atomic compare
+  if (ce < cx) { cx = ce; }
+
+#pragma omp atomic compare
+  if (cx > ce) { cx = ce; }
+
+#pragma omp atomic compare
+  if (cx < ce) { cx = ce; }
+
+#pragma omp atomic compare
+  if (cx == ce) { cx = cd; }
+
+// }
+
+// unsigned char
+// {
+
+#pragma omp atomic compare
+  ucx = uce > ucx ? uce : ucx;
+
+#pragma omp atomic compare
+  ucx = uce < ucx ? uce : ucx;
+
+#pragma omp atomic compare
+  ucx = ucx > uce ? uce : ucx;
+
+#pragma omp atomic compare
+  ucx = ucx < uce ? uce : ucx;
+
+#pragma omp atomic compare
+  ucx = ucx == uce ? ucd : ucx;
+
+#pragma omp atomic compare
+  if (uce > ucx) { ucx = uce; }
+
+#pragma omp atomic compare
+  if (uce < ucx) { ucx = uce; }
+
+#pragma omp atomic compare
+  if (ucx > uce) { ucx = uce; }
+
+#pragma omp atomic compare
+  if (ucx < uce) { ucx = uce; }
+
+#pragma omp atomic compare
+  if (ucx == uce) { ucx = ucd; }
+
+// }
+
+// short
+// {
+
+#pragma omp atomic compare
+  sx = se > sx ? se : sx;
+
+#pragma omp atomic compare
+  sx = se < sx ? se : sx;
+
+#pragma omp atomic compare
+  sx = sx > se ? se : sx;
+
+#pragma omp atomic compare
+  sx = sx < se ? se : sx;
+
+#pragma omp atomic compare
+  sx = sx == se ? sd : sx;
+
+#pragma omp atomic compare
+  if (se > sx) { sx = se; }
+
+#pragma omp atomic compare
+  if (se < sx) { sx = se; }
+
+#pragma omp atomic compare
+  if (sx > se) { sx = se; }
+
+#pragma omp atomic compare
+  if (sx < se) { sx = se; }
+
+#pragma omp atomic compare
+  if (sx == se) { sx = sd; }
+
+// }
+
+// unsigned short
+// {
+
+#pragma omp atomic compare
+  usx = use > usx ? use : usx;
+
+#pragma omp atomic compare
+  usx = use < usx ? use : usx;
+
+#pragma omp atomic compare
+  usx = usx > use ? use : usx;
+
+#pragma omp atomic compare
+  usx = usx < use ? use : usx;
+
+#pragma omp atomic compare
+  usx = usx == use ? usd : usx;
+
+#pragma omp atomic compare
+  if (use > usx) { usx = use; }
+
+#pragma omp atomic compare
+  if (use < usx) { usx = use; }
+
+#pragma omp atomic compare
+  if (usx > use) { usx = use; }
+
+#pragma omp atomic compare
+  if (usx < use) { usx = use; }
+
+#pragma omp atomic compare
+  if (usx == use) { usx = usd; }
+
+// }
+
+// int
+// {
+
+#pragma omp atomic compare
+  ix = ie > ix ? ie : ix;
+
+#pragma omp atomic compare
+  ix = ie < ix ? ie : ix;
+
+#pragma omp atomic compare
+  ix = ix > ie ? ie : ix;
+
+#pragma omp atomic compare
+  ix = ix < ie ? ie : ix;
+
+#pragma omp atomic compare
+  ix = ix == ie ? id : ix;
+
+#pragma omp atomic compare
+  if (ie > ix) { ix = ie; }
+
+#pragma omp atomic compare
+  if (ie < ix) { ix = ie; }
+
+#pragma omp atomic compare
+  if (ix > ie) { ix = ie; }
+
+#pragma omp atomic compare
+  if (ix < ie) { ix = ie; }
+
+#pragma omp atomic compare
+  if (ix == ie) { ix = id; }
+
+// }
+
+// unsigned int
+// {
+
+#pragma omp atomic compare
+  uix = uie > uix ? uie : uix;
+
+#pragma omp atomic compare
+  uix = uie < uix ? uie : uix;
+
+#pragma omp atomic compare
+  uix = uix > uie ? uie : uix;
+
+#pragma omp atomic compare
+  uix = uix < uie ? uie : uix;
+
+#pragma omp atomic compare
+  uix = uix == uie ? uid : uix;
+
+#pragma omp atomic compare
+  if (uie > uix) { uix = uie; }
+
+#pragma omp atomic compare
+  if (uie < uix) { uix = uie; }
+
+#pragma omp atomic compare
+  if (uix > uie) { uix = uie; }
+
+#pragma omp atomic compare
+  if (uix < uie) { uix = uie; }
+
+#pragma omp atomic compare
+  if (uix == uie) { uix = uid; }
+
+// }
+
+// long
+// {
+
+#pragma omp atomic compare
+  lx = le > lx ? le : lx;
+
+#pragma omp atomic compare
+  lx = le < lx ? le : lx;
+
+#pragma omp atomic compare
+  lx = lx > le ? le : lx;
+
+#pragma omp atomic compare
+  lx = lx < le ? le : lx;
+
+#pragma omp atomic compare
+  lx = lx == le ? ld : lx;
+
+#pragma omp atomic compare
+  if (le > lx) { lx = le; }
+
+#pragma omp atomic compare
+  if (le < lx) { lx = le; }
+
+#pragma omp atomic compare
+  if (lx > le) { lx = le; }
+
+#pragma omp atomic compare
+  if (lx < le) { lx = le; }
+
+#pragma omp atomic compare
+  if (lx == le) { lx = ld; }
+
+// }
+
+// unsigned long
+// {
+
+#pragma omp atomic compare
+  ulx = ule > ulx ? ule : ulx;
+
+#pragma omp atomic compare
+  ulx = ule < ulx ? ule : ulx;
+
+#pragma omp atomic compare
+  ulx = ulx > ule ? ule : ulx;
+
+#pragma omp atomic compare
+  ulx = ulx < ule ? ule : ulx;
+
+#pragma omp atomic compare
+  ulx = ulx == ule ? uld : ulx;
+
+#pragma omp atomic compare
+  if (ule > ulx) { ulx = ule; }
+
+#pragma omp atomic compare
+  if (ule < ulx) { ulx = ule; }
+
+#pragma omp atomic compare
+  if (ulx > ule) { ulx = ule; }
+
+#pragma omp atomic compare
+  if (ulx < ule) { ulx = ule; }
+
+#pragma omp atomic compare
+  if (ulx == ule) { ulx = uld; }
+
+// }
+
+// long long
+// {
+
+#pragma omp atomic compare
+  llx = lle > llx ? lle : llx;
+
+#pragma omp atomic compare
+  llx = lle < llx ? lle : llx;
+
+#pragma omp atomic compare
+  llx = llx > lle ? lle : llx;
+
+#pragma omp atomic compare
+  llx = llx < lle ? lle : llx;
+
+#pragma omp atomic compare
+  llx = llx == lle ? lld : llx;
+
+#pragma omp atomic compare
+  if (lle > llx) { llx = lle; }
+
+#pragma omp atomic compare
+  if (lle < llx) { llx = lle; }
+
+#pragma omp atomic compare
+  if (llx > lle) { llx = lle; }
+
+#pragma omp atomic compare
+  if (llx < lle) { llx = lle; }
+
+#pragma omp atomic compare
+  if (llx == lle) { llx = lld; }
+
+// }
+
+// unsigned long long
+// {
+
+#pragma omp atomic compare
+  ullx = ulle > ullx ? ulle : ullx;
+
+#pragma omp atomic compare
+  ullx = ulle < ullx ? ulle : ullx;
+
+#pragma omp atomic compare
+  ullx = ullx > ulle ? ulle : ullx;
+
+#pragma omp atomic compare
+  ullx = ullx < ulle ? ulle : ullx;
+
+#pragma omp atomic compare
+  ullx = ullx == ulle ? ulld : ullx;
+
+#pragma omp atomic compare
+  if (ulle > ullx) { ullx = ulle; }
+
+#pragma omp atomic compare
+  if (ulle < ullx) { ullx = ulle; }
+
+#pragma omp atomic compare
+  if (ullx > ulle) { ullx = ulle; }
+
+#pragma omp atomic compare
+  if (ullx < ulle) { ullx = ulle; }
+
+#pragma omp atomic compare
+  if (ullx == ulle) { ullx = ulld; }
+
+// }
+}

clang/tools/libclang/CIndex.cpp

 void OMPClauseEnqueue::VisitOMPReadClause(const OMPReadClause *) {}
 
 void OMPClauseEnqueue::VisitOMPWriteClause(const OMPWriteClause *) {}
 
 void OMPClauseEnqueue::VisitOMPUpdateClause(const OMPUpdateClause *) {}
 
 void OMPClauseEnqueue::VisitOMPCaptureClause(const OMPCaptureClause *) {}
 
+void OMPClauseEnqueue::VisitOMPCompareClause(const OMPCompareClause *) {}
+
 void OMPClauseEnqueue::VisitOMPSeqCstClause(const OMPSeqCstClause *) {}
 
 void OMPClauseEnqueue::VisitOMPAcqRelClause(const OMPAcqRelClause *) {}
 
 void OMPClauseEnqueue::VisitOMPAcquireClause(const OMPAcquireClause *) {}
 
 void OMPClauseEnqueue::VisitOMPReleaseClause(const OMPReleaseClause *) {}
 

llvm/include/llvm/Frontend/OpenMP/OMP.td

 def OMPC_Untied : Clause<"untied"> { let clangClass = "OMPUntiedClause"; }
 def OMPC_Mergeable : Clause<"mergeable"> {
   let clangClass = "OMPMergeableClause";
 }
 def OMPC_Read : Clause<"read"> { let clangClass = "OMPReadClause"; }
 def OMPC_Write : Clause<"write"> { let clangClass = "OMPWriteClause"; }
 def OMPC_Update : Clause<"update"> { let clangClass = "OMPUpdateClause"; }
 def OMPC_Capture : Clause<"capture"> { let clangClass = "OMPCaptureClause"; }
+def OMPC_Compare : Clause<"compare"> { let clangClass = "OMPCompareClause"; }
 def OMPC_SeqCst : Clause<"seq_cst"> { let clangClass = "OMPSeqCstClause"; }
 def OMPC_AcqRel : Clause<"acq_rel"> { let clangClass = "OMPAcqRelClause"; }
 def OMPC_Acquire : Clause<"acquire"> { let clangClass = "OMPAcquireClause"; }
 def OMPC_Release : Clause<"release"> { let clangClass = "OMPReleaseClause"; }
 def OMPC_Relaxed : Clause<"relaxed"> { let clangClass = "OMPRelaxedClause"; }
 def OMPC_Depend : Clause<"depend"> {
   let clangClass = "OMPDependClause";
   let flangClass = "OmpDependClause";
def OMP_Ordered : Directive<"ordered"> {
   ];
 }
 def OMP_Atomic : Directive<"atomic"> {
   let allowedClauses = [
     VersionedClause<OMPC_Read>,
     VersionedClause<OMPC_Write>,
     VersionedClause<OMPC_Update>,
     VersionedClause<OMPC_Capture>,
+    VersionedClause<OMPC_Compare, 51>
   ];
   let allowedOnceClauses = [
     VersionedClause<OMPC_SeqCst>,
     VersionedClause<OMPC_AcqRel, 50>,
     VersionedClause<OMPC_Acquire, 50>,
     VersionedClause<OMPC_Release, 50>,
     VersionedClause<OMPC_Relaxed, 50>,
     VersionedClause<OMPC_Hint, 50>