[OpenACC] Implement 'nohost' construct AST/Sema 'nohost' is only valid on routine, and states that the compiler shouldn't compile this routine for the host. It has no arguments, so no checking is required besides putting it in the AST.

diff --git a/clang/include/clang/AST/OpenACCClause.h b/clang/include/clang/AST/OpenACCClause.h
index 0fc638b73..b2cf621 100644
--- a/clang/include/clang/AST/OpenACCClause.h
+++ b/clang/include/clang/AST/OpenACCClause.h

@@ -162,6 +162,26 @@
     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 1d9e3b0..f049653 100644
--- a/clang/include/clang/Basic/OpenACCClauses.def
+++ b/clang/include/clang/Basic/OpenACCClauses.def

@@ -56,6 +56,7 @@
 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 cef241e..fd2c38a 100644
--- a/clang/lib/AST/OpenACCClause.cpp
+++ b/clang/lib/AST/OpenACCClause.cpp

@@ -534,6 +534,13 @@
   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 @@
   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 f9aa7aa..574f67f 100644
--- a/clang/lib/AST/StmtProfile.cpp
+++ b/clang/lib/AST/StmtProfile.cpp

@@ -2719,6 +2719,8 @@
     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 c07919c..91f3f14 100644
--- a/clang/lib/AST/TextNodeDumper.cpp
+++ b/clang/lib/AST/TextNodeDumper.cpp

@@ -423,6 +423,7 @@
     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 149481d..582681f 100644
--- a/clang/lib/Sema/SemaOpenACCClause.cpp
+++ b/clang/lib/Sema/SemaOpenACCClause.cpp

@@ -475,6 +475,14 @@
       return false;
     }
   }
+  case OpenACCClauseKind::NoHost: {
+    switch (DirectiveKind) {
+    case OpenACCDirectiveKind::Routine:
+      return true;
+    default:
+      return false;
+    }
+  }
   }
 
   default:
@@ -1286,6 +1294,12 @@
                                    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 79075c2..f3e3307 100644
--- a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
+++ b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp

@@ -1109,6 +1109,12 @@
                                        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 b99acf6..9591fd4 100644
--- a/clang/lib/Sema/TreeTransform.h
+++ b/clang/lib/Sema/TreeTransform.h

@@ -11875,6 +11875,11 @@
     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 2c186e4..2ac9754 100644
--- a/clang/lib/Serialization/ASTReader.cpp
+++ b/clang/lib/Serialization/ASTReader.cpp

@@ -12733,6 +12733,8 @@
   }
   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 @@
                                                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 c3e67ff..0aa115e 100644
--- a/clang/lib/Serialization/ASTWriter.cpp
+++ b/clang/lib/Serialization/ASTWriter.cpp

@@ -8783,6 +8783,7 @@
   }
   case OpenACCClauseKind::Seq:
   case OpenACCClauseKind::Independent:
+  case OpenACCClauseKind::NoHost:
   case OpenACCClauseKind::Auto:
   case OpenACCClauseKind::Finalize:
   case OpenACCClauseKind::IfPresent:
@@ -8843,7 +8844,6 @@
     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 c199db7..7397187 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 @@
 // 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 @@
 }
 // 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 @@
 #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 @@
 
 // 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 @@
 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 f59317c..8fe2c36 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 @@
 // 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 @@
 // 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 @@
 // 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 87e566bf..5390f4d 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 @@
 // 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 @@
 // 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 e536fa6..f412a38 100644
--- a/clang/tools/libclang/CIndex.cpp
+++ b/clang/tools/libclang/CIndex.cpp

@@ -2973,6 +2973,7 @@
 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(