diff --git a/clang/include/clang/AST/OpenACCClause.h b/clang/include/clang/AST/OpenACCClause.h
index 401b8e904a1b7..07587849eb121 100644
--- a/clang/include/clang/AST/OpenACCClause.h
+++ b/clang/include/clang/AST/OpenACCClause.h
@@ -145,6 +145,17 @@ class OpenACCIfClause : public OpenACCClauseWithCondition {
                                  SourceLocation EndLoc);
 };
 
+/// A 'self' clause, which has an optional condition expression.
+class OpenACCSelfClause : public OpenACCClauseWithCondition {
+  OpenACCSelfClause(SourceLocation BeginLoc, SourceLocation LParenLoc,
+                    Expr *ConditionExpr, SourceLocation EndLoc);
+
+public:
+  static OpenACCSelfClause *Create(const ASTContext &C, SourceLocation BeginLoc,
+                                   SourceLocation LParenLoc,
+                                   Expr *ConditionExpr, SourceLocation EndLoc);
+};
+
 template <class Impl> class OpenACCClauseVisitor {
   Impl &getDerived() { return static_cast<Impl &>(*this); }
 
@@ -159,53 +170,13 @@ template <class Impl> class OpenACCClauseVisitor {
       return;
 
     switch (C->getClauseKind()) {
-    case OpenACCClauseKind::Default:
-      VisitDefaultClause(*cast<OpenACCDefaultClause>(C));
-      return;
-    case OpenACCClauseKind::If:
-      VisitIfClause(*cast<OpenACCIfClause>(C));
-      return;
-    case OpenACCClauseKind::Finalize:
-    case OpenACCClauseKind::IfPresent:
-    case OpenACCClauseKind::Seq:
-    case OpenACCClauseKind::Independent:
-    case OpenACCClauseKind::Auto:
-    case OpenACCClauseKind::Worker:
-    case OpenACCClauseKind::Vector:
-    case OpenACCClauseKind::NoHost:
-    case OpenACCClauseKind::Self:
-    case OpenACCClauseKind::Copy:
-    case OpenACCClauseKind::UseDevice:
-    case OpenACCClauseKind::Attach:
-    case OpenACCClauseKind::Delete:
-    case OpenACCClauseKind::Detach:
-    case OpenACCClauseKind::Device:
-    case OpenACCClauseKind::DevicePtr:
-    case OpenACCClauseKind::DeviceResident:
-    case OpenACCClauseKind::FirstPrivate:
-    case OpenACCClauseKind::Host:
-    case OpenACCClauseKind::Link:
-    case OpenACCClauseKind::NoCreate:
-    case OpenACCClauseKind::Present:
-    case OpenACCClauseKind::Private:
-    case OpenACCClauseKind::CopyOut:
-    case OpenACCClauseKind::CopyIn:
-    case OpenACCClauseKind::Create:
-    case OpenACCClauseKind::Reduction:
-    case OpenACCClauseKind::Collapse:
-    case OpenACCClauseKind::Bind:
-    case OpenACCClauseKind::VectorLength:
-    case OpenACCClauseKind::NumGangs:
-    case OpenACCClauseKind::NumWorkers:
-    case OpenACCClauseKind::DeviceNum:
-    case OpenACCClauseKind::DefaultAsync:
-    case OpenACCClauseKind::DeviceType:
-    case OpenACCClauseKind::DType:
-    case OpenACCClauseKind::Async:
-    case OpenACCClauseKind::Tile:
-    case OpenACCClauseKind::Gang:
-    case OpenACCClauseKind::Wait:
-    case OpenACCClauseKind::Invalid:
+#define VISIT_CLAUSE(CLAUSE_NAME)                                              \
+  case OpenACCClauseKind::CLAUSE_NAME:                                         \
+    Visit##CLAUSE_NAME##Clause(*cast<OpenACC##CLAUSE_NAME##Clause>(C));        \
+    return;
+#include "clang/Basic/OpenACCClauses.def"
+
+    default:
       llvm_unreachable("Clause visitor not yet implemented");
     }
     llvm_unreachable("Invalid Clause kind");
diff --git a/clang/include/clang/AST/StmtOpenACC.h b/clang/include/clang/AST/StmtOpenACC.h
index 419cb6cada0bc..66f8f844e0b29 100644
--- a/clang/include/clang/AST/StmtOpenACC.h
+++ b/clang/include/clang/AST/StmtOpenACC.h
@@ -142,9 +142,7 @@ class OpenACCComputeConstruct final
                           Stmt *StructuredBlock)
       : OpenACCAssociatedStmtConstruct(OpenACCComputeConstructClass, K, Start,
                                        End, StructuredBlock) {
-    assert((K == OpenACCDirectiveKind::Parallel ||
-            K == OpenACCDirectiveKind::Serial ||
-            K == OpenACCDirectiveKind::Kernels) &&
+    assert(isOpenACCComputeDirectiveKind(K) &&
            "Only parallel, serial, and kernels constructs should be "
            "represented by this type");
 
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index 5ec0218aedfe8..44f802c0c28e8 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -12274,4 +12274,8 @@ def note_acc_branch_into_compute_construct
     : Note<"invalid branch into OpenACC Compute Construct">;
 def note_acc_branch_out_of_compute_construct
     : Note<"invalid branch out of OpenACC Compute Construct">;
+def warn_acc_if_self_conflict
+    : Warning<"OpenACC construct 'self' has no effect when an 'if' clause "
+              "evaluates to true">,
+      InGroup<DiagGroup<"openacc-self-if-potential-conflict">>;
 } // end of sema component.
diff --git a/clang/include/clang/Basic/OpenACCClauses.def b/clang/include/clang/Basic/OpenACCClauses.def
index 7fd2720e02ce2..378495d2c0909 100644
--- a/clang/include/clang/Basic/OpenACCClauses.def
+++ b/clang/include/clang/Basic/OpenACCClauses.def
@@ -17,5 +17,6 @@
 
 VISIT_CLAUSE(Default)
 VISIT_CLAUSE(If)
+VISIT_CLAUSE(Self)
 
 #undef VISIT_CLAUSE
diff --git a/clang/include/clang/Basic/OpenACCKinds.h b/clang/include/clang/Basic/OpenACCKinds.h
index 3414df9999170..e3f7417843328 100644
--- a/clang/include/clang/Basic/OpenACCKinds.h
+++ b/clang/include/clang/Basic/OpenACCKinds.h
@@ -146,6 +146,12 @@ inline llvm::raw_ostream &operator<<(llvm::raw_ostream &Out,
   return printOpenACCDirectiveKind(Out, K);
 }
 
+inline bool isOpenACCComputeDirectiveKind(OpenACCDirectiveKind K) {
+  return K == OpenACCDirectiveKind::Parallel ||
+         K == OpenACCDirectiveKind::Serial ||
+         K == OpenACCDirectiveKind::Kernels;
+}
+
 enum class OpenACCAtomicKind {
   Read,
   Write,
diff --git a/clang/include/clang/Sema/SemaOpenACC.h b/clang/include/clang/Sema/SemaOpenACC.h
index c1fe0f5b9c0f6..329dc3945fa2a 100644
--- a/clang/include/clang/Sema/SemaOpenACC.h
+++ b/clang/include/clang/Sema/SemaOpenACC.h
@@ -44,7 +44,8 @@ class SemaOpenACC : public SemaBase {
       Expr *ConditionExpr;
     };
 
-    std::variant<DefaultDetails, ConditionDetails> Details;
+    std::variant<std::monostate, DefaultDetails, ConditionDetails> Details =
+        std::monostate{};
 
   public:
     OpenACCParsedClause(OpenACCDirectiveKind DirKind,
@@ -72,8 +73,17 @@ class SemaOpenACC : public SemaBase {
     }
 
     Expr *getConditionExpr() {
-      assert(ClauseKind == OpenACCClauseKind::If &&
+      assert((ClauseKind == OpenACCClauseKind::If ||
+              (ClauseKind == OpenACCClauseKind::Self &&
+               DirKind != OpenACCDirectiveKind::Update)) &&
              "Parsed clause kind does not have a condition expr");
+
+      // 'self' has an optional ConditionExpr, so be tolerant of that. This will
+      // assert in variant otherwise.
+      if (ClauseKind == OpenACCClauseKind::Self &&
+          std::holds_alternative<std::monostate>(Details))
+        return nullptr;
+
       return std::get<ConditionDetails>(Details).ConditionExpr;
     }
 
@@ -87,7 +97,9 @@ class SemaOpenACC : public SemaBase {
     }
 
     void setConditionDetails(Expr *ConditionExpr) {
-      assert(ClauseKind == OpenACCClauseKind::If &&
+      assert((ClauseKind == OpenACCClauseKind::If ||
+              (ClauseKind == OpenACCClauseKind::Self &&
+               DirKind != OpenACCDirectiveKind::Update)) &&
              "Parsed clause kind does not have a condition expr");
       // In C++ we can count on this being a 'bool', but in C this gets left as
       // some sort of scalar that codegen will have to take care of converting.
diff --git a/clang/lib/AST/OpenACCClause.cpp b/clang/lib/AST/OpenACCClause.cpp
index dcb512cb51417..9c259c8f9bd0a 100644
--- a/clang/lib/AST/OpenACCClause.cpp
+++ b/clang/lib/AST/OpenACCClause.cpp
@@ -48,6 +48,26 @@ OpenACCIfClause::OpenACCIfClause(SourceLocation BeginLoc,
          "Condition expression type not scalar/dependent");
 }
 
+OpenACCSelfClause *OpenACCSelfClause::Create(const ASTContext &C,
+                                             SourceLocation BeginLoc,
+                                             SourceLocation LParenLoc,
+                                             Expr *ConditionExpr,
+                                             SourceLocation EndLoc) {
+  void *Mem = C.Allocate(sizeof(OpenACCIfClause), alignof(OpenACCIfClause));
+  return new (Mem)
+      OpenACCSelfClause(BeginLoc, LParenLoc, ConditionExpr, EndLoc);
+}
+
+OpenACCSelfClause::OpenACCSelfClause(SourceLocation BeginLoc,
+                                     SourceLocation LParenLoc,
+                                     Expr *ConditionExpr, SourceLocation EndLoc)
+    : OpenACCClauseWithCondition(OpenACCClauseKind::Self, BeginLoc, LParenLoc,
+                                 ConditionExpr, EndLoc) {
+  assert((!ConditionExpr || ConditionExpr->isInstantiationDependent() ||
+          ConditionExpr->getType()->isScalarType()) &&
+         "Condition expression type not scalar/dependent");
+}
+
 OpenACCClause::child_range OpenACCClause::children() {
   switch (getClauseKind()) {
   default:
@@ -72,3 +92,9 @@ void OpenACCClausePrinter::VisitDefaultClause(const OpenACCDefaultClause &C) {
 void OpenACCClausePrinter::VisitIfClause(const OpenACCIfClause &C) {
   OS << "if(" << C.getConditionExpr() << ")";
 }
+
+void OpenACCClausePrinter::VisitSelfClause(const OpenACCSelfClause &C) {
+  OS << "self";
+  if (const Expr *CondExpr = C.getConditionExpr())
+    OS << "(" << CondExpr << ")";
+}
diff --git a/clang/lib/AST/StmtProfile.cpp b/clang/lib/AST/StmtProfile.cpp
index 789e4634bd293..b26d804c6f079 100644
--- a/clang/lib/AST/StmtProfile.cpp
+++ b/clang/lib/AST/StmtProfile.cpp
@@ -2491,6 +2491,11 @@ void OpenACCClauseProfiler::VisitIfClause(const OpenACCIfClause &Clause) {
          "if clause requires a valid condition expr");
   Profiler.VisitStmt(Clause.getConditionExpr());
 }
+
+void OpenACCClauseProfiler::VisitSelfClause(const OpenACCSelfClause &Clause) {
+  if (Clause.hasConditionExpr())
+    Profiler.VisitStmt(Clause.getConditionExpr());
+}
 } // namespace
 
 void StmtProfiler::VisitOpenACCComputeConstruct(
diff --git a/clang/lib/AST/TextNodeDumper.cpp b/clang/lib/AST/TextNodeDumper.cpp
index 688daa64d6197..ff5b3df2d6dfa 100644
--- a/clang/lib/AST/TextNodeDumper.cpp
+++ b/clang/lib/AST/TextNodeDumper.cpp
@@ -398,6 +398,7 @@ void TextNodeDumper::Visit(const OpenACCClause *C) {
       OS << '(' << cast<OpenACCDefaultClause>(C)->getDefaultClauseKind() << ')';
       break;
     case OpenACCClauseKind::If:
+    case OpenACCClauseKind::Self:
       // The condition expression will be printed as a part of the 'children',
       // but print 'clause' here so it is clear what is happening from the dump.
       OS << " clause";
diff --git a/clang/lib/Parse/ParseOpenACC.cpp b/clang/lib/Parse/ParseOpenACC.cpp
index 91f2b8afcf0c2..123be476e928e 100644
--- a/clang/lib/Parse/ParseOpenACC.cpp
+++ b/clang/lib/Parse/ParseOpenACC.cpp
@@ -835,19 +835,23 @@ Parser::OpenACCClauseParseResult Parser::ParseOpenACCClauseParams(
     case OpenACCClauseKind::Default: {
       Token DefKindTok = getCurToken();
 
-      if (expectIdentifierOrKeyword(*this))
-        break;
+      if (expectIdentifierOrKeyword(*this)) {
+        Parens.skipToEnd();
+        return OpenACCCanContinue();
+      }
 
       ConsumeToken();
 
       OpenACCDefaultClauseKind DefKind =
           getOpenACCDefaultClauseKind(DefKindTok);
 
-      if (DefKind == OpenACCDefaultClauseKind::Invalid)
+      if (DefKind == OpenACCDefaultClauseKind::Invalid) {
         Diag(DefKindTok, diag::err_acc_invalid_default_clause_kind);
-      else
-        ParsedClause.setDefaultDetails(DefKind);
+        Parens.skipToEnd();
+        return OpenACCCanContinue();
+      }
 
+      ParsedClause.setDefaultDetails(DefKind);
       break;
     }
     case OpenACCClauseKind::If: {
@@ -977,6 +981,8 @@ Parser::OpenACCClauseParseResult Parser::ParseOpenACCClauseParams(
       case OpenACCClauseKind::Self: {
         assert(DirKind != OpenACCDirectiveKind::Update);
         ExprResult CondExpr = ParseOpenACCConditionExpr();
+        ParsedClause.setConditionDetails(CondExpr.isUsable() ? CondExpr.get()
+                                                             : nullptr);
 
         if (CondExpr.isInvalid()) {
           Parens.skipToEnd();
diff --git a/clang/lib/Sema/SemaOpenACC.cpp b/clang/lib/Sema/SemaOpenACC.cpp
index 1249136c87650..59f65eaf47a6d 100644
--- a/clang/lib/Sema/SemaOpenACC.cpp
+++ b/clang/lib/Sema/SemaOpenACC.cpp
@@ -15,6 +15,7 @@
 #include "clang/AST/StmtOpenACC.h"
 #include "clang/Basic/DiagnosticSema.h"
 #include "clang/Sema/Sema.h"
+#include "llvm/Support/Casting.h"
 
 using namespace clang;
 
@@ -76,6 +77,19 @@ bool doesClauseApplyToDirective(OpenACCDirectiveKind DirectiveKind,
     default:
       return false;
     }
+  case OpenACCClauseKind::Self:
+    switch (DirectiveKind) {
+    case OpenACCDirectiveKind::Parallel:
+    case OpenACCDirectiveKind::Serial:
+    case OpenACCDirectiveKind::Kernels:
+    case OpenACCDirectiveKind::Update:
+    case OpenACCDirectiveKind::ParallelLoop:
+    case OpenACCDirectiveKind::SerialLoop:
+    case OpenACCDirectiveKind::KernelsLoop:
+      return true;
+    default:
+      return false;
+    }
   default:
     // Do nothing so we can go to the 'unimplemented' diagnostic instead.
     return true;
@@ -121,9 +135,7 @@ SemaOpenACC::ActOnClause(ArrayRef<const OpenACCClause *> ExistingClauses,
     // Restrictions only properly implemented on 'compute' constructs, and
     // 'compute' constructs are the only construct that can do anything with
     // this yet, so skip/treat as unimplemented in this case.
-    if (Clause.getDirectiveKind() != OpenACCDirectiveKind::Parallel &&
-        Clause.getDirectiveKind() != OpenACCDirectiveKind::Serial &&
-        Clause.getDirectiveKind() != OpenACCDirectiveKind::Kernels)
+    if (!isOpenACCComputeDirectiveKind(Clause.getDirectiveKind()))
       break;
 
     // Don't add an invalid clause to the AST.
@@ -146,9 +158,7 @@ SemaOpenACC::ActOnClause(ArrayRef<const OpenACCClause *> ExistingClauses,
     // Restrictions only properly implemented on 'compute' constructs, and
     // 'compute' constructs are the only construct that can do anything with
     // this yet, so skip/treat as unimplemented in this case.
-    if (Clause.getDirectiveKind() != OpenACCDirectiveKind::Parallel &&
-        Clause.getDirectiveKind() != OpenACCDirectiveKind::Serial &&
-        Clause.getDirectiveKind() != OpenACCDirectiveKind::Kernels)
+    if (!isOpenACCComputeDirectiveKind(Clause.getDirectiveKind()))
       break;
 
     // There is no prose in the standard that says duplicates aren't allowed,
@@ -160,12 +170,54 @@ SemaOpenACC::ActOnClause(ArrayRef<const OpenACCClause *> ExistingClauses,
     // The parser has ensured that we have a proper condition expr, so there
     // isn't really much to do here.
 
-    // TODO OpenACC: When we implement 'self', this clauses causes us to
-    // 'ignore' the self clause, so we should implement a warning here.
+    // If the 'if' clause is true, it makes the 'self' clause have no effect,
+    // diagnose that here.
+    // TODO OpenACC: When we add these two to other constructs, we might not
+    // want to warn on this (for example, 'update').
+    const auto *Itr =
+        llvm::find_if(ExistingClauses, llvm::IsaPred<OpenACCSelfClause>);
+    if (Itr != ExistingClauses.end()) {
+      Diag(Clause.getBeginLoc(), diag::warn_acc_if_self_conflict);
+      Diag((*Itr)->getBeginLoc(), diag::note_acc_previous_clause_here);
+    }
+
     return OpenACCIfClause::Create(
         getASTContext(), Clause.getBeginLoc(), Clause.getLParenLoc(),
         Clause.getConditionExpr(), Clause.getEndLoc());
   }
+
+  case OpenACCClauseKind::Self: {
+    // Restrictions only properly implemented on 'compute' constructs, and
+    // 'compute' constructs are the only construct that can do anything with
+    // this yet, so skip/treat as unimplemented in this case.
+    if (!isOpenACCComputeDirectiveKind(Clause.getDirectiveKind()))
+      break;
+
+    // TODO OpenACC: When we implement this for 'update', this takes a
+    // 'var-list' instead of a condition expression, so semantics/handling has
+    // to happen differently here.
+
+    // There is no prose in the standard that says duplicates aren't allowed,
+    // but this diagnostic is present in other compilers, as well as makes
+    // sense.
+    if (checkAlreadyHasClauseOfKind(*this, ExistingClauses, Clause))
+      return nullptr;
+
+    // If the 'if' clause is true, it makes the 'self' clause have no effect,
+    // diagnose that here.
+    // TODO OpenACC: When we add these two to other constructs, we might not
+    // want to warn on this (for example, 'update').
+    const auto *Itr =
+        llvm::find_if(ExistingClauses, llvm::IsaPred<OpenACCIfClause>);
+    if (Itr != ExistingClauses.end()) {
+      Diag(Clause.getBeginLoc(), diag::warn_acc_if_self_conflict);
+      Diag((*Itr)->getBeginLoc(), diag::note_acc_previous_clause_here);
+    }
+
+    return OpenACCSelfClause::Create(
+        getASTContext(), Clause.getBeginLoc(), Clause.getLParenLoc(),
+        Clause.getConditionExpr(), Clause.getEndLoc());
+  }
   default:
     break;
   }
diff --git a/clang/lib/Sema/TreeTransform.h b/clang/lib/Sema/TreeTransform.h
index acc2d7ff9d427..0c7fdb357235e 100644
--- a/clang/lib/Sema/TreeTransform.h
+++ b/clang/lib/Sema/TreeTransform.h
@@ -11088,6 +11088,77 @@ OMPClause *TreeTransform<Derived>::TransformOMPXBareClause(OMPXBareClause *C) {
 //===----------------------------------------------------------------------===//
 // OpenACC transformation
 //===----------------------------------------------------------------------===//
+namespace {
+template <typename Derived>
+class OpenACCClauseTransform final
+    : public OpenACCClauseVisitor<OpenACCClauseTransform<Derived>> {
+  TreeTransform<Derived> &Self;
+  SemaOpenACC::OpenACCParsedClause &ParsedClause;
+  OpenACCClause *NewClause = nullptr;
+
+public:
+  OpenACCClauseTransform(TreeTransform<Derived> &Self,
+                         SemaOpenACC::OpenACCParsedClause &PC)
+      : Self(Self), ParsedClause(PC) {}
+
+  OpenACCClause *CreatedClause() const { return NewClause; }
+
+#define VISIT_CLAUSE(CLAUSE_NAME)                                              \
+  void Visit##CLAUSE_NAME##Clause(const OpenACC##CLAUSE_NAME##Clause &Clause);
+#include "clang/Basic/OpenACCClauses.def"
+};
+
+template <typename Derived>
+void OpenACCClauseTransform<Derived>::VisitDefaultClause(
+    const OpenACCDefaultClause &C) {
+  ParsedClause.setDefaultDetails(C.getDefaultClauseKind());
+
+  NewClause = OpenACCDefaultClause::Create(
+      Self.getSema().getASTContext(), ParsedClause.getDefaultClauseKind(),
+      ParsedClause.getBeginLoc(), ParsedClause.getLParenLoc(),
+      ParsedClause.getEndLoc());
+}
+
+template <typename Derived>
+void OpenACCClauseTransform<Derived>::VisitIfClause(const OpenACCIfClause &C) {
+  Expr *Cond = const_cast<Expr *>(C.getConditionExpr());
+  assert(Cond && "If constructed with invalid Condition");
+  Sema::ConditionResult Res = Self.TransformCondition(
+      Cond->getExprLoc(), /*Var=*/nullptr, Cond, Sema::ConditionKind::Boolean);
+
+  if (Res.isInvalid() || !Res.get().second)
+    return;
+
+  ParsedClause.setConditionDetails(Res.get().second);
+
+  NewClause = OpenACCIfClause::Create(
+      Self.getSema().getASTContext(), ParsedClause.getBeginLoc(),
+      ParsedClause.getLParenLoc(), ParsedClause.getConditionExpr(),
+      ParsedClause.getEndLoc());
+}
+
+template <typename Derived>
+void OpenACCClauseTransform<Derived>::VisitSelfClause(
+    const OpenACCSelfClause &C) {
+
+  if (C.hasConditionExpr()) {
+    Expr *Cond = const_cast<Expr *>(C.getConditionExpr());
+    Sema::ConditionResult Res =
+        Self.TransformCondition(Cond->getExprLoc(), /*Var=*/nullptr, Cond,
+                                Sema::ConditionKind::Boolean);
+
+    if (Res.isInvalid() || !Res.get().second)
+      return;
+
+    ParsedClause.setConditionDetails(Res.get().second);
+  }
+
+  NewClause = OpenACCSelfClause::Create(
+      Self.getSema().getASTContext(), ParsedClause.getBeginLoc(),
+      ParsedClause.getLParenLoc(), ParsedClause.getConditionExpr(),
+      ParsedClause.getEndLoc());
+}
+} // namespace
 template <typename Derived>
 OpenACCClause *TreeTransform<Derived>::TransformOpenACCClause(
     ArrayRef<const OpenACCClause *> ExistingClauses,
@@ -11100,33 +11171,10 @@ OpenACCClause *TreeTransform<Derived>::TransformOpenACCClause(
   if (const auto *WithParms = dyn_cast<OpenACCClauseWithParams>(OldClause))
     ParsedClause.setLParenLoc(WithParms->getLParenLoc());
 
-  switch (OldClause->getClauseKind()) {
-  case OpenACCClauseKind::Default:
-    // There is nothing to do here as nothing dependent can appear in this
-    // clause. So just set the values so Sema can set the right value.
-    ParsedClause.setDefaultDetails(
-        cast<OpenACCDefaultClause>(OldClause)->getDefaultClauseKind());
-    break;
-  case OpenACCClauseKind::If: {
-    Expr *Cond = const_cast<Expr *>(
-        cast<OpenACCIfClause>(OldClause)->getConditionExpr());
-    assert(Cond && "If constructed with invalid Condition");
-    Sema::ConditionResult Res =
-        TransformCondition(Cond->getExprLoc(), /*Var=*/nullptr, Cond,
-                           Sema::ConditionKind::Boolean);
-
-    if (Res.isInvalid() || !Res.get().second)
-      return nullptr;
-
-    ParsedClause.setConditionDetails(Res.get().second);
-    break;
-  }
-  default:
-    assert(false && "Unhandled OpenACC clause in TreeTransform");
-    return nullptr;
-  }
+  OpenACCClauseTransform<Derived> Transform{*this, ParsedClause};
+  Transform.Visit(OldClause);
 
-  return getSema().OpenACC().ActOnClause(ExistingClauses, ParsedClause);
+  return Transform.CreatedClause();
 }
 
 template <typename Derived>
diff --git a/clang/lib/Serialization/ASTReader.cpp b/clang/lib/Serialization/ASTReader.cpp
index f47d540ea4b86..cf0726460bfca 100644
--- a/clang/lib/Serialization/ASTReader.cpp
+++ b/clang/lib/Serialization/ASTReader.cpp
@@ -11781,6 +11781,12 @@ OpenACCClause *ASTRecordReader::readOpenACCClause() {
     return OpenACCIfClause::Create(getContext(), BeginLoc, LParenLoc, CondExpr,
                                    EndLoc);
   }
+  case OpenACCClauseKind::Self: {
+    SourceLocation LParenLoc = readSourceLocation();
+    Expr *CondExpr = readBool() ? readSubExpr() : nullptr;
+    return OpenACCSelfClause::Create(getContext(), BeginLoc, LParenLoc,
+                                     CondExpr, EndLoc);
+  }
   case OpenACCClauseKind::Finalize:
   case OpenACCClauseKind::IfPresent:
   case OpenACCClauseKind::Seq:
@@ -11789,7 +11795,6 @@ OpenACCClause *ASTRecordReader::readOpenACCClause() {
   case OpenACCClauseKind::Worker:
   case OpenACCClauseKind::Vector:
   case OpenACCClauseKind::NoHost:
-  case OpenACCClauseKind::Self:
   case OpenACCClauseKind::Copy:
   case OpenACCClauseKind::UseDevice:
   case OpenACCClauseKind::Attach:
diff --git a/clang/lib/Serialization/ASTWriter.cpp b/clang/lib/Serialization/ASTWriter.cpp
index ce6fa1feb1eeb..b2a078b6d80f4 100644
--- a/clang/lib/Serialization/ASTWriter.cpp
+++ b/clang/lib/Serialization/ASTWriter.cpp
@@ -7524,6 +7524,14 @@ void ASTRecordWriter::writeOpenACCClause(const OpenACCClause *C) {
     AddStmt(const_cast<Expr *>(IC->getConditionExpr()));
     return;
   }
+  case OpenACCClauseKind::Self: {
+    const auto *SC = cast<OpenACCIfClause>(C);
+    writeSourceLocation(SC->getLParenLoc());
+    writeBool(SC->hasConditionExpr());
+    if (SC->hasConditionExpr())
+      AddStmt(const_cast<Expr *>(SC->getConditionExpr()));
+    return;
+  }
   case OpenACCClauseKind::Finalize:
   case OpenACCClauseKind::IfPresent:
   case OpenACCClauseKind::Seq:
@@ -7532,7 +7540,6 @@ void ASTRecordWriter::writeOpenACCClause(const OpenACCClause *C) {
   case OpenACCClauseKind::Worker:
   case OpenACCClauseKind::Vector:
   case OpenACCClauseKind::NoHost:
-  case OpenACCClauseKind::Self:
   case OpenACCClauseKind::Copy:
   case OpenACCClauseKind::UseDevice:
   case OpenACCClauseKind::Attach:
diff --git a/clang/test/ParserOpenACC/parse-clauses.c b/clang/test/ParserOpenACC/parse-clauses.c
index 2369df58308a7..4462f0df540f2 100644
--- a/clang/test/ParserOpenACC/parse-clauses.c
+++ b/clang/test/ParserOpenACC/parse-clauses.c
@@ -376,16 +376,13 @@ void SelfClause() {
 #pragma acc serial self(i > j, seq
   for(;;){}
 
-  // expected-warning@+2{{left operand of comma operator has no effect}}
-  // expected-warning@+1{{OpenACC clause 'self' not yet implemented, clause ignored}}
+  // expected-warning@+1{{left operand of comma operator has no effect}}
 #pragma acc serial self(i, j)
   for(;;){}
 
-  // expected-warning@+1{{OpenACC clause 'self' not yet implemented, clause ignored}}
 #pragma acc serial self(i > j)
   for(;;){}
 
-  // expected-warning@+2{{OpenACC clause 'self' not yet implemented, clause ignored}}
   // expected-warning@+1{{OpenACC clause 'seq' not yet implemented, clause ignored}}
 #pragma acc serial self(1+5>3), seq
   for(;;){}
diff --git a/clang/test/SemaOpenACC/compute-construct-clause-ast.cpp b/clang/test/SemaOpenACC/compute-construct-clause-ast.cpp
index 018f0b68c7810..6d2efcf81eb6e 100644
--- a/clang/test/SemaOpenACC/compute-construct-clause-ast.cpp
+++ b/clang/test/SemaOpenACC/compute-construct-clause-ast.cpp
@@ -110,6 +110,50 @@ void TemplFunc() {
   // CHECK-NEXT: CXXBoolLiteralExpr
   // CHECK-NEXT: NullStmt
 
+#pragma acc serial self
+  while(true);
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}}serial
+  // CHECK-NEXT: self clause
+  // CHECK-NEXT: WhileStmt
+  // CHECK-NEXT: CXXBoolLiteralExpr
+  // CHECK-NEXT: NullStmt
+
+#pragma acc kernels self(T::SomeFloat)
+  while(true);
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}}kernels
+  // CHECK-NEXT: self clause
+  // CHECK-NEXT: DependentScopeDeclRefExpr{{.*}} '<dependent type>' lvalue
+  // CHECK-NEXT: NestedNameSpecifier TypeSpec 'T'
+  // CHECK-NEXT: WhileStmt
+  // CHECK-NEXT: CXXBoolLiteralExpr
+  // CHECK-NEXT: NullStmt
+
+#pragma acc parallel self(T::SomeFloat) if (T::SomeFloat)
+  while(true);
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}}parallel
+  // CHECK-NEXT: self clause
+  // CHECK-NEXT: DependentScopeDeclRefExpr{{.*}} '<dependent type>' lvalue
+  // CHECK-NEXT: NestedNameSpecifier TypeSpec 'T'
+  // CHECK-NEXT: if clause
+  // CHECK-NEXT: DependentScopeDeclRefExpr{{.*}} '<dependent type>' lvalue
+  // CHECK-NEXT: NestedNameSpecifier TypeSpec 'T'
+  // CHECK-NEXT: WhileStmt
+  // CHECK-NEXT: CXXBoolLiteralExpr
+  // CHECK-NEXT: NullStmt
+
+#pragma acc serial if(T::SomeFloat) self(T::SomeFloat)
+  while(true);
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}}serial
+  // CHECK-NEXT: if clause
+  // CHECK-NEXT: DependentScopeDeclRefExpr{{.*}} '<dependent type>' lvalue
+  // CHECK-NEXT: NestedNameSpecifier TypeSpec 'T'
+  // CHECK-NEXT: self clause
+  // CHECK-NEXT: DependentScopeDeclRefExpr{{.*}} '<dependent type>' lvalue
+  // CHECK-NEXT: NestedNameSpecifier TypeSpec 'T'
+  // CHECK-NEXT: WhileStmt
+  // CHECK-NEXT: CXXBoolLiteralExpr
+  // CHECK-NEXT: NullStmt
+
   // Match the instantiation:
   // CHECK: FunctionDecl{{.*}}TemplFunc{{.*}}implicit_instantiation
   // CHECK-NEXT: TemplateArgument type 'InstTy'
@@ -171,6 +215,53 @@ void TemplFunc() {
   // CHECK-NEXT: WhileStmt
   // CHECK-NEXT: CXXBoolLiteralExpr
   // CHECK-NEXT: NullStmt
+
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}}serial
+  // CHECK-NEXT: self clause
+  // CHECK-NEXT: WhileStmt
+  // CHECK-NEXT: CXXBoolLiteralExpr
+  // CHECK-NEXT: NullStmt
+
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}}kernels
+  // CHECK-NEXT: self clause
+  // CHECK-NEXT: ImplicitCastExpr{{.*}}'bool' <FloatingToBoolean>
+  // CHECK-NEXT: ImplicitCastExpr{{.*}}'float' <LValueToRValue>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'const float' lvalue Var{{.*}} 'SomeFloat' 'const float'
+  // CHECK-NEXT: NestedNameSpecifier TypeSpec 'InstTy'
+  // CHECK-NEXT: WhileStmt
+  // CHECK-NEXT: CXXBoolLiteralExpr
+  // CHECK-NEXT: NullStmt
+
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}}parallel
+  // CHECK-NEXT: self clause
+  // CHECK-NEXT: ImplicitCastExpr{{.*}}'bool' <FloatingToBoolean>
+  // CHECK-NEXT: ImplicitCastExpr{{.*}}'float' <LValueToRValue>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'const float' lvalue Var{{.*}} 'SomeFloat' 'const float'
+  // CHECK-NEXT: NestedNameSpecifier TypeSpec 'InstTy'
+  // CHECK-NEXT: if clause
+  // CHECK-NEXT: ImplicitCastExpr{{.*}}'bool' <FloatingToBoolean>
+  // CHECK-NEXT: ImplicitCastExpr{{.*}}'float' <LValueToRValue>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'const float' lvalue Var{{.*}} 'SomeFloat' 'const float'
+  // CHECK-NEXT: NestedNameSpecifier TypeSpec 'InstTy'
+  // CHECK-NEXT: WhileStmt
+  // CHECK-NEXT: CXXBoolLiteralExpr
+  // CHECK-NEXT: NullStmt
+
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}}serial
+  // CHECK-NEXT: if clause
+  // CHECK-NEXT: ImplicitCastExpr{{.*}}'bool' <FloatingToBoolean>
+  // CHECK-NEXT: ImplicitCastExpr{{.*}}'float' <LValueToRValue>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'const float' lvalue Var{{.*}} 'SomeFloat' 'const float'
+  // CHECK-NEXT: NestedNameSpecifier TypeSpec 'InstTy'
+  // CHECK-NEXT: self clause
+  // CHECK-NEXT: ImplicitCastExpr{{.*}}'bool' <FloatingToBoolean>
+  // CHECK-NEXT: ImplicitCastExpr{{.*}}'float' <LValueToRValue>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'const float' lvalue Var{{.*}} 'SomeFloat' 'const float'
+  // CHECK-NEXT: NestedNameSpecifier TypeSpec 'InstTy'
+  // CHECK-NEXT: WhileStmt
+  // CHECK-NEXT: CXXBoolLiteralExpr
+  // CHECK-NEXT: NullStmt
+
 }
 
 struct BoolConversion{ operator bool() const;};
diff --git a/clang/test/SemaOpenACC/compute-construct-self-clause.c b/clang/test/SemaOpenACC/compute-construct-self-clause.c
new file mode 100644
index 0000000000000..fbed2953419a2
--- /dev/null
+++ b/clang/test/SemaOpenACC/compute-construct-self-clause.c
@@ -0,0 +1,82 @@
+// RUN: %clang_cc1 %s -fopenacc -verify
+
+void BoolExpr(int *I, float *F) {
+  typedef struct {} SomeStruct;
+  struct C{};
+  // expected-error@+1{{expected expression}}
+#pragma acc parallel self (struct C f())
+  while(0);
+
+  // expected-error@+1{{unexpected type name 'SomeStruct': expected expression}}
+#pragma acc serial self (SomeStruct)
+  while(0);
+
+  // expected-error@+1{{unexpected type name 'SomeStruct': expected expression}}
+#pragma acc serial self (SomeStruct())
+  while(0);
+
+  SomeStruct S;
+  // expected-error@+1{{statement requires expression of scalar type ('SomeStruct' invalid)}}
+#pragma acc serial self (S)
+  while(0);
+
+#pragma acc parallel self (I)
+  while(0);
+
+#pragma acc serial self (F)
+  while(0);
+
+#pragma acc kernels self (*I < *F)
+  while(0);
+}
+
+void WarnMaybeNotUsed(int val1, int val2) {
+
+  // expected-warning@+2{{OpenACC construct 'self' has no effect when an 'if' clause evaluates to true}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc parallel self if(val1)
+  while(0);
+
+  // expected-warning@+2{{OpenACC construct 'self' has no effect when an 'if' clause evaluates to true}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc parallel self(val1) if(val1)
+  while(0);
+
+  // expected-warning@+2{{OpenACC construct 'self' has no effect when an 'if' clause evaluates to true}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc parallel if(val1) self
+  while(0);
+
+  // expected-warning@+2{{OpenACC construct 'self' has no effect when an 'if' clause evaluates to true}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc parallel if(val1) self(val2)
+  while(0);
+
+  // The below don't warn because one side or the other has an error, thus is
+  // not added to the AST.
+
+  // expected-error@+1{{use of undeclared identifier 'invalid'}}
+#pragma acc parallel self if(invalid)
+  while(0);
+
+  // expected-error@+1{{use of undeclared identifier 'invalid'}}
+#pragma acc parallel self(invalid) if(val1)
+  while(0);
+
+  // expected-error@+2{{expected expression}}
+  // expected-error@+1{{use of undeclared identifier 'invalid'}}
+#pragma acc parallel self() if(invalid)
+  while(0);
+
+  // expected-error@+1{{use of undeclared identifier 'invalid'}}
+#pragma acc parallel if(invalid) self
+  while(0);
+
+  // expected-error@+1{{use of undeclared identifier 'invalid'}}
+#pragma acc parallel if(val2) self(invalid)
+  while(0);
+
+  // expected-error@+1{{use of undeclared identifier 'invalid'}}
+#pragma acc parallel if(invalid) self(val1)
+  while(0);
+}
diff --git a/clang/test/SemaOpenACC/compute-construct-self-clause.cpp b/clang/test/SemaOpenACC/compute-construct-self-clause.cpp
new file mode 100644
index 0000000000000..60edbdc2b1191
--- /dev/null
+++ b/clang/test/SemaOpenACC/compute-construct-self-clause.cpp
@@ -0,0 +1,99 @@
+// RUN: %clang_cc1 %s -fopenacc -verify
+
+struct NoBoolConversion{};
+struct BoolConversion{
+  operator bool();
+};
+
+template <typename T, typename U>
+void BoolExpr() {
+  // expected-error@+1{{value of type 'NoBoolConversion' is not contextually convertible to 'bool'}}
+#pragma acc parallel self (NoBoolConversion{})
+  while(0);
+  // expected-error@+2{{no member named 'NotValid' in 'NoBoolConversion'}}
+  // expected-note@#INST{{in instantiation of function template specialization}}
+#pragma acc parallel self (T::NotValid)
+  while(0);
+
+#pragma acc parallel self (BoolConversion{})
+  while(0);
+
+  // expected-error@+1{{value of type 'NoBoolConversion' is not contextually convertible to 'bool'}}
+#pragma acc parallel self (T{})
+  while(0);
+
+#pragma acc parallel self (U{})
+  while(0);
+}
+
+struct HasBool {
+  static constexpr bool B = true;
+};
+
+template<typename T>
+void WarnMaybeNotUsed() {
+  // expected-warning@+2{{OpenACC construct 'self' has no effect when an 'if' clause evaluates to true}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc parallel self if(T::B)
+  while(0);
+
+  // expected-warning@+2{{OpenACC construct 'self' has no effect when an 'if' clause evaluates to true}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc parallel self(T::B) if(T::B)
+  while(0);
+
+  // expected-warning@+2{{OpenACC construct 'self' has no effect when an 'if' clause evaluates to true}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc parallel if(T::B) self
+  while(0);
+
+  // expected-warning@+2{{OpenACC construct 'self' has no effect when an 'if' clause evaluates to true}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc parallel if(T::B) self(T::B)
+  while(0);
+
+  // We still warn in the cases of dependent failures, since the diagnostic
+  // happens immediately rather than during instantiation.
+
+  // expected-error@+4{{no member named 'Invalid' in 'HasBool'}}
+  // expected-note@#NOT_USED_INST{{in instantiation of function template specialization 'WarnMaybeNotUsed<HasBool>' requested here}}
+  // expected-warning@+2{{OpenACC construct 'self' has no effect when an 'if' clause evaluates to true}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc parallel self if(T::Invalid)
+  while(0);
+
+  // expected-error@+3{{no member named 'Invalid' in 'HasBool'}}
+  // expected-warning@+2{{OpenACC construct 'self' has no effect when an 'if' clause evaluates to true}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc parallel self(T::Invalid) if(T::B)
+  while(0);
+
+  // expected-error@+3{{no member named 'Invalid' in 'HasBool'}}
+  // expected-warning@+2{{OpenACC construct 'self' has no effect when an 'if' clause evaluates to true}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc parallel self(T::B) if(T::Invalid)
+  while(0);
+
+  // expected-error@+3{{no member named 'Invalid' in 'HasBool'}}
+  // expected-warning@+2{{OpenACC construct 'self' has no effect when an 'if' clause evaluates to true}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc parallel if(T::Invalid) self
+  while(0);
+
+  // expected-error@+3{{no member named 'Invalid' in 'HasBool'}}
+  // expected-warning@+2{{OpenACC construct 'self' has no effect when an 'if' clause evaluates to true}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc parallel if(T::Invalid) self(T::B)
+  while(0);
+
+  // expected-error@+3{{no member named 'Invalid' in 'HasBool'}}
+  // expected-warning@+2{{OpenACC construct 'self' has no effect when an 'if' clause evaluates to true}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc parallel if(T::B) self(T::Invalid)
+  while(0);
+}
+
+void Instantiate() {
+  BoolExpr<NoBoolConversion, BoolConversion>(); // #INST
+  WarnMaybeNotUsed<HasBool>(); // #NOT_USED_INST
+}
diff --git a/clang/tools/libclang/CIndex.cpp b/clang/tools/libclang/CIndex.cpp
index f304786ff9dff..2ef599d2cd26f 100644
--- a/clang/tools/libclang/CIndex.cpp
+++ b/clang/tools/libclang/CIndex.cpp
@@ -2791,6 +2791,10 @@ void OpenACCClauseEnqueue::VisitDefaultClause(const OpenACCDefaultClause &C) {}
 void OpenACCClauseEnqueue::VisitIfClause(const OpenACCIfClause &C) {
   Visitor.AddStmt(C.getConditionExpr());
 }
+void OpenACCClauseEnqueue::VisitSelfClause(const OpenACCSelfClause &C) {
+  if (C.hasConditionExpr())
+    Visitor.AddStmt(C.getConditionExpr());
+}
 } // namespace
 
 void EnqueueVisitor::EnqueueChildren(const OpenACCClause *C) {