diff --git a/clang/include/clang/AST/OpenACCClause.h b/clang/include/clang/AST/OpenACCClause.h
index 0fc638b73bbd..b2cf621bc0a7 100644
--- a/clang/include/clang/AST/OpenACCClause.h
+++ b/clang/include/clang/AST/OpenACCClause.h
@@ -162,6 +162,26 @@ public:
     return const_child_range(const_child_iterator(), const_child_iterator());
   }
 };
+// Represents the 'nohost' clause.
+class OpenACCNoHostClause : public OpenACCClause {
+protected:
+  OpenACCNoHostClause(SourceLocation BeginLoc, SourceLocation EndLoc)
+      : OpenACCClause(OpenACCClauseKind::NoHost, BeginLoc, EndLoc) {}
+
+public:
+  static bool classof(const OpenACCClause *C) {
+    return C->getClauseKind() == OpenACCClauseKind::NoHost;
+  }
+  static OpenACCNoHostClause *
+  Create(const ASTContext &Ctx, SourceLocation BeginLoc, SourceLocation EndLoc);
+
+  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());
+  }
+};
 
 /// Represents a clause that has a list of parameters.
 class OpenACCClauseWithParams : public OpenACCClause {
diff --git a/clang/include/clang/Basic/OpenACCClauses.def b/clang/include/clang/Basic/OpenACCClauses.def
index 1d9e3b09a749..f04965363f25 100644
--- a/clang/include/clang/Basic/OpenACCClauses.def
+++ b/clang/include/clang/Basic/OpenACCClauses.def
@@ -56,6 +56,7 @@ VISIT_CLAUSE(IfPresent)
 VISIT_CLAUSE(Independent)
 VISIT_CLAUSE(Link)
 VISIT_CLAUSE(NoCreate)
+VISIT_CLAUSE(NoHost)
 VISIT_CLAUSE(NumGangs)
 VISIT_CLAUSE(NumWorkers)
 VISIT_CLAUSE(Present)
diff --git a/clang/lib/AST/OpenACCClause.cpp b/clang/lib/AST/OpenACCClause.cpp
index cef241eface9..fd2c38a0e64e 100644
--- a/clang/lib/AST/OpenACCClause.cpp
+++ b/clang/lib/AST/OpenACCClause.cpp
@@ -534,6 +534,13 @@ OpenACCSeqClause *OpenACCSeqClause::Create(const ASTContext &C,
   return new (Mem) OpenACCSeqClause(BeginLoc, EndLoc);
 }
 
+OpenACCNoHostClause *OpenACCNoHostClause::Create(const ASTContext &C,
+                                                 SourceLocation BeginLoc,
+                                                 SourceLocation EndLoc) {
+  void *Mem = C.Allocate(sizeof(OpenACCNoHostClause));
+  return new (Mem) OpenACCNoHostClause(BeginLoc, EndLoc);
+}
+
 OpenACCGangClause *
 OpenACCGangClause::Create(const ASTContext &C, SourceLocation BeginLoc,
                           SourceLocation LParenLoc,
@@ -871,6 +878,10 @@ void OpenACCClausePrinter::VisitSeqClause(const OpenACCSeqClause &C) {
   OS << "seq";
 }
 
+void OpenACCClausePrinter::VisitNoHostClause(const OpenACCNoHostClause &C) {
+  OS << "nohost";
+}
+
 void OpenACCClausePrinter::VisitCollapseClause(const OpenACCCollapseClause &C) {
   OS << "collapse(";
   if (C.hasForce())
diff --git a/clang/lib/AST/StmtProfile.cpp b/clang/lib/AST/StmtProfile.cpp
index f9aa7aa079e6..574f67f4274e 100644
--- a/clang/lib/AST/StmtProfile.cpp
+++ b/clang/lib/AST/StmtProfile.cpp
@@ -2719,6 +2719,8 @@ void OpenACCClauseProfiler::VisitIndependentClause(
     const OpenACCIndependentClause &Clause) {}
 
 void OpenACCClauseProfiler::VisitSeqClause(const OpenACCSeqClause &Clause) {}
+void OpenACCClauseProfiler::VisitNoHostClause(
+    const OpenACCNoHostClause &Clause) {}
 
 void OpenACCClauseProfiler::VisitGangClause(const OpenACCGangClause &Clause) {
   for (unsigned I = 0; I < Clause.getNumExprs(); ++I) {
diff --git a/clang/lib/AST/TextNodeDumper.cpp b/clang/lib/AST/TextNodeDumper.cpp
index c07919c4b3e3..91f3f14c6b45 100644
--- a/clang/lib/AST/TextNodeDumper.cpp
+++ b/clang/lib/AST/TextNodeDumper.cpp
@@ -423,6 +423,7 @@ void TextNodeDumper::Visit(const OpenACCClause *C) {
     case OpenACCClauseKind::FirstPrivate:
     case OpenACCClauseKind::Link:
     case OpenACCClauseKind::NoCreate:
+    case OpenACCClauseKind::NoHost:
     case OpenACCClauseKind::NumGangs:
     case OpenACCClauseKind::NumWorkers:
     case OpenACCClauseKind::Present:
diff --git a/clang/lib/Sema/SemaOpenACCClause.cpp b/clang/lib/Sema/SemaOpenACCClause.cpp
index 149481d337b1..582681f247b3 100644
--- a/clang/lib/Sema/SemaOpenACCClause.cpp
+++ b/clang/lib/Sema/SemaOpenACCClause.cpp
@@ -475,6 +475,14 @@ bool doesClauseApplyToDirective(OpenACCDirectiveKind DirectiveKind,
       return false;
     }
   }
+  case OpenACCClauseKind::NoHost: {
+    switch (DirectiveKind) {
+    case OpenACCDirectiveKind::Routine:
+      return true;
+    default:
+      return false;
+    }
+  }
   }
 
   default:
@@ -1286,6 +1294,12 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitAutoClause(
                                    Clause.getEndLoc());
 }
 
+OpenACCClause *SemaOpenACCClauseVisitor::VisitNoHostClause(
+    SemaOpenACC::OpenACCParsedClause &Clause) {
+  return OpenACCNoHostClause::Create(Ctx, Clause.getBeginLoc(),
+                                     Clause.getEndLoc());
+}
+
 OpenACCClause *SemaOpenACCClauseVisitor::VisitIndependentClause(
     SemaOpenACC::OpenACCParsedClause &Clause) {
   // OpenACC 3.3 2.9:
diff --git a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
index 79075c23c3ef..f3e33078c55c 100644
--- a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
+++ b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
@@ -1109,6 +1109,12 @@ void OpenACCDeclClauseInstantiator::VisitSeqClause(const OpenACCSeqClause &C) {
                                        ParsedClause.getBeginLoc(),
                                        ParsedClause.getEndLoc());
 }
+void OpenACCDeclClauseInstantiator::VisitNoHostClause(
+    const OpenACCNoHostClause &C) {
+  NewClause = OpenACCNoHostClause::Create(SemaRef.getASTContext(),
+                                          ParsedClause.getBeginLoc(),
+                                          ParsedClause.getEndLoc());
+}
 
 void OpenACCDeclClauseInstantiator::VisitWorkerClause(
     const OpenACCWorkerClause &C) {
diff --git a/clang/lib/Sema/TreeTransform.h b/clang/lib/Sema/TreeTransform.h
index b99acf6b723f..9591fd4cfcc1 100644
--- a/clang/lib/Sema/TreeTransform.h
+++ b/clang/lib/Sema/TreeTransform.h
@@ -11875,6 +11875,11 @@ void OpenACCClauseTransform<Derived>::VisitDeviceResidentClause(
     const OpenACCDeviceResidentClause &C) {
   llvm_unreachable("device_resident clause not valid unless a decl transform");
 }
+template <typename Derived>
+void OpenACCClauseTransform<Derived>::VisitNoHostClause(
+    const OpenACCNoHostClause &C) {
+  llvm_unreachable("device_resident clause not valid unless a decl transform");
+}
 
 template <typename Derived>
 void OpenACCClauseTransform<Derived>::VisitCopyInClause(
diff --git a/clang/lib/Serialization/ASTReader.cpp b/clang/lib/Serialization/ASTReader.cpp
index 2c186e4047dd..2ac9754f02ee 100644
--- a/clang/lib/Serialization/ASTReader.cpp
+++ b/clang/lib/Serialization/ASTReader.cpp
@@ -12733,6 +12733,8 @@ OpenACCClause *ASTRecordReader::readOpenACCClause() {
   }
   case OpenACCClauseKind::Seq:
     return OpenACCSeqClause::Create(getContext(), BeginLoc, EndLoc);
+  case OpenACCClauseKind::NoHost:
+    return OpenACCNoHostClause::Create(getContext(), BeginLoc, EndLoc);
   case OpenACCClauseKind::Finalize:
     return OpenACCFinalizeClause::Create(getContext(), BeginLoc, EndLoc);
   case OpenACCClauseKind::IfPresent:
@@ -12795,7 +12797,6 @@ OpenACCClause *ASTRecordReader::readOpenACCClause() {
                                                LParenLoc, VarList, EndLoc);
   }
 
-  case OpenACCClauseKind::NoHost:
   case OpenACCClauseKind::Bind:
   case OpenACCClauseKind::Invalid:
     llvm_unreachable("Clause serialization not yet implemented");
diff --git a/clang/lib/Serialization/ASTWriter.cpp b/clang/lib/Serialization/ASTWriter.cpp
index c3e67ffa50ce..0aa115ecadf8 100644
--- a/clang/lib/Serialization/ASTWriter.cpp
+++ b/clang/lib/Serialization/ASTWriter.cpp
@@ -8783,6 +8783,7 @@ void ASTRecordWriter::writeOpenACCClause(const OpenACCClause *C) {
   }
   case OpenACCClauseKind::Seq:
   case OpenACCClauseKind::Independent:
+  case OpenACCClauseKind::NoHost:
   case OpenACCClauseKind::Auto:
   case OpenACCClauseKind::Finalize:
   case OpenACCClauseKind::IfPresent:
@@ -8843,7 +8844,6 @@ void ASTRecordWriter::writeOpenACCClause(const OpenACCClause *C) {
     return;
   }
 
-  case OpenACCClauseKind::NoHost:
   case OpenACCClauseKind::Bind:
   case OpenACCClauseKind::Invalid:
     llvm_unreachable("Clause serialization not yet implemented");
diff --git a/clang/test/AST/ast-print-openacc-routine-construct.cpp b/clang/test/AST/ast-print-openacc-routine-construct.cpp
index c199db71a4e4..739718724f17 100644
--- a/clang/test/AST/ast-print-openacc-routine-construct.cpp
+++ b/clang/test/AST/ast-print-openacc-routine-construct.cpp
@@ -4,8 +4,8 @@ auto Lambda = [](){};
 // CHECK: #pragma acc routine(Lambda) worker
 #pragma acc routine(Lambda) worker
 int function();
-// CHECK: #pragma acc routine(function) vector
-#pragma acc routine (function) vector
+// CHECK: #pragma acc routine(function) vector nohost
+#pragma acc routine (function) vector nohost
 
 namespace NS {
   int NSFunc();
@@ -13,8 +13,8 @@ auto Lambda = [](){};
 }
 // CHECK: #pragma acc routine(NS::NSFunc) seq
 #pragma acc routine(NS::NSFunc) seq
-// CHECK: #pragma acc routine(NS::Lambda) gang
-#pragma acc routine(NS::Lambda) gang
+// CHECK: #pragma acc routine(NS::Lambda) nohost gang
+#pragma acc routine(NS::Lambda) nohost gang
 
 constexpr int getInt() { return 1; }
 
@@ -33,8 +33,8 @@ struct S {
 #pragma acc routine(MemFunc) gang(dim:1)
 // CHECK: #pragma acc routine(StaticMemFunc) gang(dim: getInt())
 #pragma acc routine(StaticMemFunc) gang(dim:getInt())
-// CHECK: #pragma acc routine(Lambda) worker
-#pragma acc routine(Lambda) worker
+// CHECK: #pragma acc routine(Lambda) nohost worker
+#pragma acc routine(Lambda) nohost worker
 };
 
 // CHECK: #pragma acc routine(S::MemFunc) gang(dim: 1)
@@ -66,10 +66,10 @@ struct DepS {
 
 // CHECK: #pragma acc routine(DepS<T>::Lambda) vector
 #pragma acc routine(DepS<T>::Lambda) vector
-// CHECK: #pragma acc routine(DepS<T>::MemFunc) seq
-#pragma acc routine(DepS<T>::MemFunc) seq
-// CHECK: #pragma acc routine(DepS<T>::StaticMemFunc) worker
-#pragma acc routine(DepS<T>::StaticMemFunc) worker
+// CHECK: #pragma acc routine(DepS<T>::MemFunc) seq nohost
+#pragma acc routine(DepS<T>::MemFunc) seq nohost
+// CHECK: #pragma acc routine(DepS<T>::StaticMemFunc) nohost worker
+#pragma acc routine(DepS<T>::StaticMemFunc) nohost worker
 };
 
 // CHECK: #pragma acc routine(DepS<int>::Lambda) gang
@@ -84,8 +84,8 @@ template<typename T>
 void TemplFunc() {
 // CHECK: #pragma acc routine(T::MemFunc) gang(dim: T::SomethingElse())
 #pragma acc routine(T::MemFunc) gang(dim:T::SomethingElse())
-// CHECK: #pragma acc routine(T::StaticMemFunc) worker
-#pragma acc routine(T::StaticMemFunc) worker
-// CHECK: #pragma acc routine(T::Lambda) seq
-#pragma acc routine(T::Lambda) seq
+// CHECK: #pragma acc routine(T::StaticMemFunc) worker nohost
+#pragma acc routine(T::StaticMemFunc) worker nohost
+// CHECK: #pragma acc routine(T::Lambda) nohost seq
+#pragma acc routine(T::Lambda) nohost seq
 }
diff --git a/clang/test/SemaOpenACC/routine-construct-ast.cpp b/clang/test/SemaOpenACC/routine-construct-ast.cpp
index f59317cf9d93..8fe2c3695826 100644
--- a/clang/test/SemaOpenACC/routine-construct-ast.cpp
+++ b/clang/test/SemaOpenACC/routine-construct-ast.cpp
@@ -7,14 +7,16 @@
 #ifndef PCH_HELPER
 #define PCH_HELPER
 auto Lambda = [](){};
-#pragma acc routine(Lambda) worker
+#pragma acc routine(Lambda) worker nohost
 // CHECK: OpenACCRoutineDecl{{.*}} routine name_specified
 // CHECK-NEXT: DeclRefExpr{{.*}} 'Lambda' '(lambda at
 // CHECK-NEXT: worker clause
+// CHECK-NEXT: nohost clause
 int function();
-#pragma acc routine (function) vector
+#pragma acc routine (function) nohost vector
 // CHECK: OpenACCRoutineDecl{{.*}} routine name_specified
 // CHECK-NEXT: DeclRefExpr{{.*}} 'function' 'int ()'
+// CHECK-NEXT: nohost clause
 // CHECK-NEXT: vector clause
 
 namespace NS {
@@ -227,18 +229,20 @@ void TemplFunc() {
 // CHECK-NEXT: NestedNameSpecifier{{.*}} 'T'
 // CHECK-NEXT: gang clause
 // CHECK-NEXT: CallExpr{{.*}}'<dependent type>'
-#pragma acc routine(T::StaticMemFunc) worker
+#pragma acc routine(T::StaticMemFunc) nohost worker
 // CHECK-NEXT: DeclStmt
 // CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
 // CHECK-NEXT: DependentScopeDeclRefExpr{{.*}}'<dependent type>'
 // CHECK-NEXT: NestedNameSpecifier{{.*}} 'T'
+// CHECK-NEXT: nohost clause
 // CHECK-NEXT: worker clause
-#pragma acc routine(T::Lambda) seq
+#pragma acc routine(T::Lambda) seq nohost
 // CHECK-NEXT: DeclStmt
 // CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
 // CHECK-NEXT: DependentScopeDeclRefExpr{{.*}}'<dependent type>'
 // CHECK-NEXT: NestedNameSpecifier{{.*}} 'T'
 // CHECK-NEXT: seq clause
+// CHECK-NEXT: nohost clause
 
 // Instantiation:
 // CHECK: FunctionDecl{{.*}} TemplFunc 'void ()' implicit_instantiation
@@ -254,6 +258,7 @@ void TemplFunc() {
 // CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
 // CHECK-NEXT: DeclRefExpr{{.*}} 'StaticMemFunc' 'void ()'
 // CHECK-NEXT: NestedNameSpecifier{{.*}} 'S'
+// CHECK-NEXT: nohost clause
 // CHECK-NEXT: worker clause
 
 // CHECK-NEXT: DeclStmt
@@ -261,6 +266,7 @@ void TemplFunc() {
 // CHECK-NEXT: DeclRefExpr{{.*}} 'Lambda' 'const S::(lambda at
 // CHECK-NEXT: NestedNameSpecifier{{.*}} 'S'
 // CHECK-NEXT: seq clause
+// CHECK-NEXT: nohost clause
 }
 
 void usage() {
diff --git a/clang/test/SemaOpenACC/routine-construct-clauses.cpp b/clang/test/SemaOpenACC/routine-construct-clauses.cpp
index 87e566bfe81c..5390f4d32691 100644
--- a/clang/test/SemaOpenACC/routine-construct-clauses.cpp
+++ b/clang/test/SemaOpenACC/routine-construct-clauses.cpp
@@ -3,8 +3,8 @@
 void Func();
 
 #pragma acc routine(Func) worker
-#pragma acc routine(Func) vector
-#pragma acc routine(Func) seq
+#pragma acc routine(Func) vector nohost
+#pragma acc routine(Func) nohost seq
 #pragma acc routine(Func) gang
 
 // Only 1 of worker, vector, seq, gang.
@@ -56,6 +56,10 @@ void Func();
 // expected-error@+2{{OpenACC clause 'gang' may not appear on the same construct as a 'gang' clause on a 'routine' construct}}
 // expected-note@+1{{previous clause is here}}
 #pragma acc routine(Func) gang gang
+// expected-error@+1{{REQUIRED}}
+#pragma acc routine(Func)
+// expected-error@+1{{REQUIRED}}
+#pragma acc routine(Func) nohost
 
 // only the 'dim' syntax for gang is legal.
 #pragma acc routine(Func) gang(dim:1)
@@ -106,8 +110,8 @@ struct DependentT {
 // expected-error@+1{{argument to 'gang' clause dimension must be 1, 2, or 3: evaluated to -5}}
 #pragma acc routine(Func) gang(dim:T::Neg())
 // expected-error@+1{{argument to 'gang' clause dimension must be 1, 2, or 3: evaluated to 0}}
-#pragma acc routine(Func) gang(dim:T::Zero())
-#pragma acc routine(Func) gang(dim:T::One())
+#pragma acc routine(Func) gang(dim:T::Zero()) nohost
+#pragma acc routine(Func) nohost gang(dim:T::One())
 #pragma acc routine(Func) gang(dim:T::Two())
 #pragma acc routine(Func) gang(dim:T::Three())
 // expected-error@+1{{argument to 'gang' clause dimension must be 1, 2, or 3: evaluated to 4}}
diff --git a/clang/tools/libclang/CIndex.cpp b/clang/tools/libclang/CIndex.cpp
index e536fa6aaff1..f412a38a9233 100644
--- a/clang/tools/libclang/CIndex.cpp
+++ b/clang/tools/libclang/CIndex.cpp
@@ -2973,6 +2973,7 @@ void OpenACCClauseEnqueue::VisitAutoClause(const OpenACCAutoClause &C) {}
 void OpenACCClauseEnqueue::VisitIndependentClause(
     const OpenACCIndependentClause &C) {}
 void OpenACCClauseEnqueue::VisitSeqClause(const OpenACCSeqClause &C) {}
+void OpenACCClauseEnqueue::VisitNoHostClause(const OpenACCNoHostClause &C) {}
 void OpenACCClauseEnqueue::VisitFinalizeClause(const OpenACCFinalizeClause &C) {
 }
 void OpenACCClauseEnqueue::VisitIfPresentClause(