From a61f6afbee370785cf091fe46e2e022748528307 Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <thomas@codesourcery.com>
Date: Wed, 21 Jul 2021 18:30:00 +0200
Subject: [PATCH] OpenACC 'nohost' clause

Do not "compile a version of this procedure for the host".

	gcc/
	* tree-core.h (omp_clause_code): Add 'OMP_CLAUSE_NOHOST'.
	* tree.c (omp_clause_num_ops, omp_clause_code_name, walk_tree_1):
	Handle it.
	* tree-pretty-print.c (dump_omp_clause): Likewise.
	* omp-general.c (oacc_verify_routine_clauses): Likewise.
	* gimplify.c (gimplify_scan_omp_clauses)
	(gimplify_adjust_omp_clauses): Likewise.
	* tree-nested.c (convert_nonlocal_omp_clauses)
	(convert_local_omp_clauses): Likewise.
	* omp-low.c (scan_sharing_clauses): Likewise.
	* omp-offload.c (execute_oacc_device_lower): Update.
	gcc/c-family/
	* c-pragma.h (pragma_omp_clause): Add 'PRAGMA_OACC_CLAUSE_NOHOST'.
	gcc/c/
	* c-parser.c (c_parser_omp_clause_name): Handle 'nohost'.
	(c_parser_oacc_all_clauses): Handle 'PRAGMA_OACC_CLAUSE_NOHOST'.
	(OACC_ROUTINE_CLAUSE_MASK): Add 'PRAGMA_OACC_CLAUSE_NOHOST'.
	* c-typeck.c (c_finish_omp_clauses): Handle 'OMP_CLAUSE_NOHOST'.
	gcc/cp/
	* parser.c (cp_parser_omp_clause_name): Handle 'nohost'.
	(cp_parser_oacc_all_clauses): Handle 'PRAGMA_OACC_CLAUSE_NOHOST'.
	(OACC_ROUTINE_CLAUSE_MASK): Add 'PRAGMA_OACC_CLAUSE_NOHOST'.
	* pt.c (tsubst_omp_clauses): Handle 'OMP_CLAUSE_NOHOST'.
	* semantics.c (finish_omp_clauses): Likewise.
	gcc/fortran/
	* dump-parse-tree.c (show_attr): Update.
	* gfortran.h (symbol_attribute): Add 'oacc_routine_nohost' member.
	(gfc_omp_clauses): Add 'nohost' member.
	* module.c (ab_attribute): Add 'AB_OACC_ROUTINE_NOHOST'.
	(attr_bits, mio_symbol_attribute): Update.
	* openmp.c (omp_mask2): Add 'OMP_CLAUSE_NOHOST'.
	(gfc_match_omp_clauses): Handle 'OMP_CLAUSE_NOHOST'.
	(OACC_ROUTINE_CLAUSES): Add 'OMP_CLAUSE_NOHOST'.
	(gfc_match_oacc_routine): Update.
	* trans-decl.c (add_attributes_to_decl): Update.
	* trans-openmp.c (gfc_trans_omp_clauses): Likewise.
	gcc/testsuite/
	* c-c++-common/goacc/classify-routine-nohost.c: New file.
	* c-c++-common/goacc/classify-routine.c: Update.
	* c-c++-common/goacc/routine-2.c: Likewise.
	* c-c++-common/goacc/routine-nohost-1.c: New file.
	* c-c++-common/goacc/routine-nohost-2.c: Likewise.
	* g++.dg/goacc/template.C: Update.
	* gfortran.dg/goacc/classify-routine-nohost.f95: New file.
	* gfortran.dg/goacc/classify-routine.f95: Update.
	* gfortran.dg/goacc/pure-elemental-procedures-2.f90: Likewise.
	* gfortran.dg/goacc/routine-6.f90: Likewise.
	* gfortran.dg/goacc/routine-intrinsic-2.f: Likewise.
	* gfortran.dg/goacc/routine-module-1.f90: Likewise.
	* gfortran.dg/goacc/routine-module-2.f90: Likewise.
	* gfortran.dg/goacc/routine-module-3.f90: Likewise.
	* gfortran.dg/goacc/routine-module-mod-1.f90: Likewise.
	* gfortran.dg/goacc/routine-multiple-directives-1.f90: Likewise.
	* gfortran.dg/goacc/routine-multiple-directives-2.f90: Likewise.
	libgomp/
	* testsuite/libgomp.oacc-c-c++-common/routine-nohost-1.c: New
	file.
	* testsuite/libgomp.oacc-c-c++-common/routine-nohost-2.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/routine-nohost-2_2.c:
	Likewise.
	* testsuite/libgomp.oacc-fortran/routine-nohost-1.f90: Likewise.

Co-Authored-By: Joseph Myers <joseph@codesourcery.com>
Co-Authored-By: Cesar Philippidis <cesar@codesourcery.com>
---
 gcc/c-family/c-pragma.h                       |   1 +
 gcc/c/c-parser.c                              |  10 +-
 gcc/c/c-typeck.c                              |   1 +
 gcc/cp/parser.c                               |  11 +-
 gcc/cp/pt.c                                   |   1 +
 gcc/cp/semantics.c                            |   1 +
 gcc/fortran/dump-parse-tree.c                 |   2 +
 gcc/fortran/gfortran.h                        |   2 +
 gcc/fortran/module.c                          |   7 +
 gcc/fortran/openmp.c                          |  30 +++-
 gcc/fortran/trans-decl.c                      |   8 +
 gcc/fortran/trans-openmp.c                    |   2 +
 gcc/gimplify.c                                |   2 +
 gcc/omp-general.c                             |  17 ++
 gcc/omp-low.c                                 |   2 +
 gcc/omp-offload.c                             |  36 +++++
 .../goacc/classify-routine-nohost.c           |  41 +++++
 .../c-c++-common/goacc/classify-routine.c     |  10 +-
 gcc/testsuite/c-c++-common/goacc/routine-2.c  |   4 +
 .../c-c++-common/goacc/routine-nohost-1.c     |  50 ++++++
 .../c-c++-common/goacc/routine-nohost-2.c     |  96 ++++++++++++
 gcc/testsuite/g++.dg/goacc/template.C         |  15 +-
 .../goacc/classify-routine-nohost.f95         |  39 +++++
 .../gfortran.dg/goacc/classify-routine.f95    |   7 +
 .../goacc/pure-elemental-procedures-2.f90     |  24 +++
 gcc/testsuite/gfortran.dg/goacc/routine-6.f90 |  10 ++
 .../gfortran.dg/goacc/routine-intrinsic-2.f   |  10 ++
 .../gfortran.dg/goacc/routine-module-1.f90    |  14 ++
 .../gfortran.dg/goacc/routine-module-2.f90    |   6 +
 .../gfortran.dg/goacc/routine-module-3.f90    |  43 ++++-
 .../goacc/routine-module-mod-1.f90            |  60 +++++++
 .../goacc/routine-multiple-directives-1.f90   |  64 ++++++++
 .../goacc/routine-multiple-directives-2.f90   | 147 ++++++++++++++++++
 gcc/tree-core.h                               |   5 +-
 gcc/tree-nested.c                             |   6 +
 gcc/tree-pretty-print.c                       |   3 +
 gcc/tree.c                                    |   3 +
 .../routine-nohost-1.c                        |  63 ++++++++
 .../routine-nohost-2.c                        |  39 +++++
 .../routine-nohost-2_2.c                      |  18 +++
 .../libgomp.oacc-fortran/routine-nohost-1.f90 |  63 ++++++++
 41 files changed, 962 insertions(+), 11 deletions(-)
 create mode 100644 gcc/testsuite/c-c++-common/goacc/classify-routine-nohost.c
 create mode 100644 gcc/testsuite/c-c++-common/goacc/routine-nohost-1.c
 create mode 100644 gcc/testsuite/c-c++-common/goacc/routine-nohost-2.c
 create mode 100644 gcc/testsuite/gfortran.dg/goacc/classify-routine-nohost.f95
 create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/routine-nohost-1.c
 create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/routine-nohost-2.c
 create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/routine-nohost-2_2.c
 create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/routine-nohost-1.f90

diff --git a/gcc/c-family/c-pragma.h b/gcc/c-family/c-pragma.h
index e4fd3c9b740..c5d11ce0a52 100644
--- a/gcc/c-family/c-pragma.h
+++ b/gcc/c-family/c-pragma.h
@@ -160,6 +160,7 @@ enum pragma_omp_clause {
   PRAGMA_OACC_CLAUSE_HOST,
   PRAGMA_OACC_CLAUSE_INDEPENDENT,
   PRAGMA_OACC_CLAUSE_NO_CREATE,
+  PRAGMA_OACC_CLAUSE_NOHOST,
   PRAGMA_OACC_CLAUSE_NUM_GANGS,
   PRAGMA_OACC_CLAUSE_NUM_WORKERS,
   PRAGMA_OACC_CLAUSE_PRESENT,
diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index 9a56e0c04c6..92d22d1af4d 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -12744,6 +12744,8 @@ c_parser_omp_clause_name (c_parser *parser)
 	    result = PRAGMA_OACC_CLAUSE_NO_CREATE;
 	  else if (!strcmp ("nogroup", p))
 	    result = PRAGMA_OMP_CLAUSE_NOGROUP;
+	  else if (!strcmp ("nohost", p))
+	    result = PRAGMA_OACC_CLAUSE_NOHOST;
 	  else if (!strcmp ("nontemporal", p))
 	    result = PRAGMA_OMP_CLAUSE_NONTEMPORAL;
 	  else if (!strcmp ("notinbranch", p))
@@ -16248,6 +16250,11 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
 	  clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
 	  c_name = "no_create";
 	  break;
+	case PRAGMA_OACC_CLAUSE_NOHOST:
+	  clauses = c_parser_oacc_simple_clause (here, OMP_CLAUSE_NOHOST,
+						 clauses);
+	  c_name = "nohost";
+	  break;
 	case PRAGMA_OACC_CLAUSE_NUM_GANGS:
 	  clauses = c_parser_oacc_single_int_clause (parser,
 						     OMP_CLAUSE_NUM_GANGS,
@@ -17179,7 +17186,8 @@ c_parser_oacc_compute (location_t loc, c_parser *parser,
 	( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_GANG)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WORKER)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR)		\
-	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_SEQ) )
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_SEQ)			\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NOHOST) )
 
 /* Parse an OpenACC routine directive.  For named directives, we apply
    immediately to the named function.  For unnamed ones we then parse
diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c
index 4f7ed675746..5d6565bdaa9 100644
--- a/gcc/c/c-typeck.c
+++ b/gcc/c/c-typeck.c
@@ -15168,6 +15168,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	case OMP_CLAUSE_TILE:
 	case OMP_CLAUSE_IF_PRESENT:
 	case OMP_CLAUSE_FINALIZE:
+	case OMP_CLAUSE_NOHOST:
 	  pc = &OMP_CLAUSE_CHAIN (c);
 	  continue;
 
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index 821ce1771a4..45216f0a222 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -35656,6 +35656,8 @@ cp_parser_omp_clause_name (cp_parser *parser)
 	    result = PRAGMA_OACC_CLAUSE_NO_CREATE;
 	  else if (!strcmp ("nogroup", p))
 	    result = PRAGMA_OMP_CLAUSE_NOGROUP;
+	  else if (!strcmp ("nohost", p))
+	    result = PRAGMA_OACC_CLAUSE_NOHOST;
 	  else if (!strcmp ("nontemporal", p))
 	    result = PRAGMA_OMP_CLAUSE_NONTEMPORAL;
 	  else if (!strcmp ("notinbranch", p))
@@ -38879,6 +38881,11 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
 	  clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
 	  c_name = "no_create";
 	  break;
+	case PRAGMA_OACC_CLAUSE_NOHOST:
+	  clauses = cp_parser_oacc_simple_clause (here, OMP_CLAUSE_NOHOST,
+						  clauses);
+	  c_name = "nohost";
+	  break;
 	case PRAGMA_OACC_CLAUSE_NUM_GANGS:
 	  code = OMP_CLAUSE_NUM_GANGS;
 	  c_name = "num_gangs";
@@ -44866,8 +44873,8 @@ cp_parser_omp_taskloop (cp_parser *parser, cp_token *pragma_tok,
 	( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_GANG)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WORKER)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR)		\
-	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_SEQ))
-
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_SEQ)			\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NOHOST) )
 
 /* Parse the OpenACC routine pragma.  This has an optional '( name )'
    component, which must resolve to a declared namespace-scope
diff --git a/gcc/cp/pt.c b/gcc/cp/pt.c
index 94ca3bc633e..b396ddd0089 100644
--- a/gcc/cp/pt.c
+++ b/gcc/cp/pt.c
@@ -17479,6 +17479,7 @@ tsubst_omp_clauses (tree clauses, enum c_omp_region_type ort,
 	case OMP_CLAUSE_SEQ:
 	case OMP_CLAUSE_IF_PRESENT:
 	case OMP_CLAUSE_FINALIZE:
+	case OMP_CLAUSE_NOHOST:
 	  break;
 	default:
 	  gcc_unreachable ();
diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c
index 331daf81bb7..f64b084963c 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -8267,6 +8267,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	case OMP_CLAUSE_SEQ:
 	case OMP_CLAUSE_IF_PRESENT:
 	case OMP_CLAUSE_FINALIZE:
+	case OMP_CLAUSE_NOHOST:
 	  break;
 
 	case OMP_CLAUSE_MERGEABLE:
diff --git a/gcc/fortran/dump-parse-tree.c b/gcc/fortran/dump-parse-tree.c
index 26841eefb7d..8e4a101b2ae 100644
--- a/gcc/fortran/dump-parse-tree.c
+++ b/gcc/fortran/dump-parse-tree.c
@@ -926,6 +926,8 @@ show_attr (symbol_attribute *attr, const char * module)
     fputs (" ALWAYS-EXPLICIT", dumpfile);
   if (attr->is_main_program)
     fputs (" IS-MAIN-PROGRAM", dumpfile);
+  if (attr->oacc_routine_nohost)
+    fputs (" OACC-ROUTINE-NOHOST", dumpfile);
 
   /* FIXME: Still missing are oacc_routine_lop and ext_attr.  */
   fputc (')', dumpfile);
diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h
index f4a50d74f14..921aed93dc3 100644
--- a/gcc/fortran/gfortran.h
+++ b/gcc/fortran/gfortran.h
@@ -947,6 +947,7 @@ typedef struct
 
   /* OpenACC 'routine' directive's level of parallelism.  */
   ENUM_BITFIELD (oacc_routine_lop) oacc_routine_lop:3;
+  unsigned oacc_routine_nohost:1;
 
   /* Attributes set by compiler extensions (!GCC$ ATTRIBUTES).  */
   unsigned ext_attr:EXT_ATTR_NUM;
@@ -1488,6 +1489,7 @@ typedef struct gfc_omp_clauses
   unsigned async:1, gang:1, worker:1, vector:1, seq:1, independent:1;
   unsigned par_auto:1, gang_static:1;
   unsigned if_present:1, finalize:1;
+  unsigned nohost:1;
   locus loc;
 }
 gfc_omp_clauses;
diff --git a/gcc/fortran/module.c b/gcc/fortran/module.c
index 321d3256eba..1804066bc8c 100644
--- a/gcc/fortran/module.c
+++ b/gcc/fortran/module.c
@@ -2088,6 +2088,7 @@ enum ab_attribute
   AB_PDT_TEMPLATE, AB_PDT_ARRAY, AB_PDT_STRING,
   AB_OACC_ROUTINE_LOP_GANG, AB_OACC_ROUTINE_LOP_WORKER,
   AB_OACC_ROUTINE_LOP_VECTOR, AB_OACC_ROUTINE_LOP_SEQ,
+  AB_OACC_ROUTINE_NOHOST,
   AB_OMP_REQ_REVERSE_OFFLOAD, AB_OMP_REQ_UNIFIED_ADDRESS,
   AB_OMP_REQ_UNIFIED_SHARED_MEMORY, AB_OMP_REQ_DYNAMIC_ALLOCATORS,
   AB_OMP_REQ_MEM_ORDER_SEQ_CST, AB_OMP_REQ_MEM_ORDER_ACQ_REL,
@@ -2166,6 +2167,7 @@ static const mstring attr_bits[] =
     minit ("OACC_ROUTINE_LOP_WORKER", AB_OACC_ROUTINE_LOP_WORKER),
     minit ("OACC_ROUTINE_LOP_VECTOR", AB_OACC_ROUTINE_LOP_VECTOR),
     minit ("OACC_ROUTINE_LOP_SEQ", AB_OACC_ROUTINE_LOP_SEQ),
+    minit ("OACC_ROUTINE_NOHOST", AB_OACC_ROUTINE_NOHOST),
     minit ("OMP_REQ_REVERSE_OFFLOAD", AB_OMP_REQ_REVERSE_OFFLOAD),
     minit ("OMP_REQ_UNIFIED_ADDRESS", AB_OMP_REQ_UNIFIED_ADDRESS),
     minit ("OMP_REQ_UNIFIED_SHARED_MEMORY", AB_OMP_REQ_UNIFIED_SHARED_MEMORY),
@@ -2420,6 +2422,8 @@ mio_symbol_attribute (symbol_attribute *attr)
 	default:
 	  gcc_unreachable ();
 	}
+      if (attr->oacc_routine_nohost)
+	MIO_NAME (ab_attribute) (AB_OACC_ROUTINE_NOHOST, attr_bits);
 
       if (attr->flavor == FL_MODULE && gfc_current_ns->omp_requires)
 	{
@@ -2682,6 +2686,9 @@ mio_symbol_attribute (symbol_attribute *attr)
 	      verify_OACC_ROUTINE_LOP_NONE (attr->oacc_routine_lop);
 	      attr->oacc_routine_lop = OACC_ROUTINE_LOP_SEQ;
 	      break;
+	    case AB_OACC_ROUTINE_NOHOST:
+	      attr->oacc_routine_nohost = 1;
+	      break;
 	    case AB_OMP_REQ_REVERSE_OFFLOAD:
 	      gfc_omp_requires_add_clause (OMP_REQ_REVERSE_OFFLOAD,
 					   "reverse_offload",
diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
index 357a1e15e01..520a435e181 100644
--- a/gcc/fortran/openmp.c
+++ b/gcc/fortran/openmp.c
@@ -880,6 +880,7 @@ enum omp_mask2
   OMP_CLAUSE_IF_PRESENT,
   OMP_CLAUSE_FINALIZE,
   OMP_CLAUSE_ATTACH,
+  OMP_CLAUSE_NOHOST,
   /* This must come last.  */
   OMP_MASK2_LAST
 };
@@ -2083,6 +2084,13 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
 	      c->nogroup = needs_space = true;
 	      continue;
 	    }
+	  if ((mask & OMP_CLAUSE_NOHOST)
+	      && !c->nohost
+	      && gfc_match ("nohost") == MATCH_YES)
+	    {
+	      c->nohost = needs_space = true;
+	      continue;
+	    }
 	  if ((mask & OMP_CLAUSE_NOTEMPORAL)
 	      && gfc_match_omp_variable_list ("nontemporal (",
 					      &c->lists[OMP_LIST_NONTEMPORAL],
@@ -2607,7 +2615,8 @@ end:
   omp_mask (OMP_CLAUSE_ASYNC)
 #define OACC_ROUTINE_CLAUSES \
   (omp_mask (OMP_CLAUSE_GANG) | OMP_CLAUSE_WORKER | OMP_CLAUSE_VECTOR	      \
-   | OMP_CLAUSE_SEQ)
+   | OMP_CLAUSE_SEQ							      \
+   | OMP_CLAUSE_NOHOST)
 
 
 static match
@@ -2936,6 +2945,7 @@ gfc_match_oacc_routine (void)
   gfc_omp_clauses *c = NULL;
   gfc_oacc_routine_name *n = NULL;
   oacc_routine_lop lop = OACC_ROUTINE_LOP_NONE;
+  bool nohost;
 
   old_loc = gfc_current_locus;
 
@@ -3012,6 +3022,7 @@ gfc_match_oacc_routine (void)
       gfc_error ("Multiple loop axes specified for routine at %C");
       goto cleanup;
     }
+  nohost = c ? c->nohost : false;
 
   if (isym != NULL)
     {
@@ -3024,6 +3035,13 @@ gfc_match_oacc_routine (void)
 		     " clause");
 	  goto cleanup;
 	}
+      /* ..., and no 'nohost' clause.  */
+      if (nohost)
+	{
+	  gfc_error ("Intrinsic symbol specified in !$ACC ROUTINE ( NAME )"
+		     " at %C marked with incompatible NOHOST clause");
+	  goto cleanup;
+	}
     }
   else if (sym != NULL)
     {
@@ -3037,7 +3055,9 @@ gfc_match_oacc_routine (void)
 	if (n_p->sym == sym)
 	  {
 	    add = false;
-	    if (lop != gfc_oacc_routine_lop (n_p->clauses))
+	    bool nohost_p = n_p->clauses ? n_p->clauses->nohost : false;
+	    if (lop != gfc_oacc_routine_lop (n_p->clauses)
+		|| nohost != nohost_p)
 	      {
 		gfc_error ("!$ACC ROUTINE already applied at %C");
 		goto cleanup;
@@ -3047,6 +3067,7 @@ gfc_match_oacc_routine (void)
       if (add)
 	{
 	  sym->attr.oacc_routine_lop = lop;
+	  sym->attr.oacc_routine_nohost = nohost;
 
 	  n = gfc_get_oacc_routine_name ();
 	  n->sym = sym;
@@ -3061,8 +3082,10 @@ gfc_match_oacc_routine (void)
       /* For a repeated OpenACC 'routine' directive, diagnose if it doesn't
 	 match the first one.  */
       oacc_routine_lop lop_p = gfc_current_ns->proc_name->attr.oacc_routine_lop;
+      bool nohost_p = gfc_current_ns->proc_name->attr.oacc_routine_nohost;
       if (lop_p != OACC_ROUTINE_LOP_NONE
-	  && lop != lop_p)
+	  && (lop != lop_p
+	      || nohost != nohost_p))
 	{
 	  gfc_error ("!$ACC ROUTINE already applied at %C");
 	  goto cleanup;
@@ -3073,6 +3096,7 @@ gfc_match_oacc_routine (void)
 				       &old_loc))
 	goto cleanup;
       gfc_current_ns->proc_name->attr.oacc_routine_lop = lop;
+      gfc_current_ns->proc_name->attr.oacc_routine_nohost = nohost;
     }
   else
     /* Something has gone wrong, possibly a syntax error.  */
diff --git a/gcc/fortran/trans-decl.c b/gcc/fortran/trans-decl.c
index a73ce8a3f40..bf8783a35f8 100644
--- a/gcc/fortran/trans-decl.c
+++ b/gcc/fortran/trans-decl.c
@@ -1473,6 +1473,14 @@ add_attributes_to_decl (symbol_attribute sym_attr, tree list)
       tree dims = oacc_build_routine_dims (clauses);
       list = oacc_replace_fn_attrib_attr (list, dims);
     }
+
+  if (sym_attr.oacc_routine_nohost)
+    {
+      tree c = build_omp_clause (UNKNOWN_LOCATION, OMP_CLAUSE_NOHOST);
+      OMP_CLAUSE_CHAIN (c) = clauses;
+      clauses = c;
+    }
+
   if (sym_attr.omp_device_type != OMP_DEVICE_TYPE_UNSET)
     {
       tree c = build_omp_clause (UNKNOWN_LOCATION, OMP_CLAUSE_DEVICE_TYPE);
diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c
index ace4faf038a..ac3f5f35bc1 100644
--- a/gcc/fortran/trans-openmp.c
+++ b/gcc/fortran/trans-openmp.c
@@ -4297,6 +4297,8 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 	  gcc_unreachable ();
 	}
     }
+  /* OpenACC 'nohost' clauses cannot appear here.  */
+  gcc_checking_assert (!clauses->nohost);
 
   return nreverse (omp_clauses);
 }
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 5d43f76f002..21ff32ee4aa 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -10310,6 +10310,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	  }
 	  break;
 
+	case OMP_CLAUSE_NOHOST:
 	default:
 	  gcc_unreachable ();
 	}
@@ -11247,6 +11248,7 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
 	case OMP_CLAUSE_EXCLUSIVE:
 	  break;
 
+	case OMP_CLAUSE_NOHOST:
 	default:
 	  gcc_unreachable ();
 	}
diff --git a/gcc/omp-general.c b/gcc/omp-general.c
index a1bb9d8d25d..b46a537e281 100644
--- a/gcc/omp-general.c
+++ b/gcc/omp-general.c
@@ -2576,6 +2576,7 @@ oacc_verify_routine_clauses (tree fndecl, tree *clauses, location_t loc,
 			     const char *routine_str)
 {
   tree c_level = NULL_TREE;
+  tree c_nohost = NULL_TREE;
   tree c_p = NULL_TREE;
   for (tree c = *clauses; c; c_p = c, c = OMP_CLAUSE_CHAIN (c))
     switch (OMP_CLAUSE_CODE (c))
@@ -2608,6 +2609,10 @@ oacc_verify_routine_clauses (tree fndecl, tree *clauses, location_t loc,
 	    c = c_p;
 	  }
 	break;
+      case OMP_CLAUSE_NOHOST:
+	/* Don't worry about duplicate clauses here.  */
+	c_nohost = c;
+	break;
       default:
 	gcc_unreachable ();
       }
@@ -2642,6 +2647,7 @@ oacc_verify_routine_clauses (tree fndecl, tree *clauses, location_t loc,
 	 this one for compatibility.  */
       /* Collect previous directive's clauses.  */
       tree c_level_p = NULL_TREE;
+      tree c_nohost_p = NULL_TREE;
       for (tree c = TREE_VALUE (attr); c; c = OMP_CLAUSE_CHAIN (c))
 	switch (OMP_CLAUSE_CODE (c))
 	  {
@@ -2652,6 +2658,10 @@ oacc_verify_routine_clauses (tree fndecl, tree *clauses, location_t loc,
 	    gcc_checking_assert (c_level_p == NULL_TREE);
 	    c_level_p = c;
 	    break;
+	  case OMP_CLAUSE_NOHOST:
+	    gcc_checking_assert (c_nohost_p == NULL_TREE);
+	    c_nohost_p = c;
+	    break;
 	  default:
 	    gcc_unreachable ();
 	  }
@@ -2667,6 +2677,13 @@ oacc_verify_routine_clauses (tree fndecl, tree *clauses, location_t loc,
 	  c_diag_p = c_level_p;
 	  goto incompatible;
 	}
+      /* Matching 'nohost' clauses?  */
+      if ((c_nohost == NULL_TREE) != (c_nohost_p == NULL_TREE))
+	{
+	  c_diag = c_nohost;
+	  c_diag_p = c_nohost_p;
+	  goto incompatible;
+	}
       /* Compatible.  */
       return 1;
 
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index e7049c825a4..2f735bcde9c 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -1683,6 +1683,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
 	  break;
 
 	case OMP_CLAUSE__CACHE_:
+	case OMP_CLAUSE_NOHOST:
 	default:
 	  gcc_unreachable ();
 	}
@@ -1869,6 +1870,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
 	  break;
 
 	case OMP_CLAUSE__CACHE_:
+	case OMP_CLAUSE_NOHOST:
 	default:
 	  gcc_unreachable ();
 	}
diff --git a/gcc/omp-offload.c b/gcc/omp-offload.c
index 0320ea6ab85..bfbb0112e24 100644
--- a/gcc/omp-offload.c
+++ b/gcc/omp-offload.c
@@ -1981,6 +1981,42 @@ execute_oacc_device_lower ()
 	gcc_unreachable ();
     }
 
+  if (is_oacc_routine)
+    {
+      tree attr = lookup_attribute ("omp declare target",
+				    DECL_ATTRIBUTES (current_function_decl));
+      gcc_checking_assert (attr);
+      tree clauses = TREE_VALUE (attr);
+      gcc_checking_assert (clauses);
+
+      /* Should this OpenACC routine be discarded?  */
+      bool discard = false;
+
+      tree clause_nohost = omp_find_clause (clauses, OMP_CLAUSE_NOHOST);
+      if (dump_file)
+	fprintf (dump_file,
+		 "OpenACC routine '%s' %s '%s' clause.\n",
+		 lang_hooks.decl_printable_name (current_function_decl, 2),
+		 clause_nohost ? "has" : "doesn't have",
+		 omp_clause_code_name[OMP_CLAUSE_NOHOST]);
+      /* Host compiler, 'nohost' clause?  */
+#ifndef ACCEL_COMPILER
+      if (clause_nohost)
+	discard = true;
+#endif
+
+      if (dump_file)
+	fprintf (dump_file,
+		 "OpenACC routine '%s' %sdiscarded.\n",
+		 lang_hooks.decl_printable_name (current_function_decl, 2),
+		 discard ? "" : "not ");
+      if (discard)
+	{
+	  TREE_ASM_WRITTEN (current_function_decl) = 1;
+	  return TODO_discard_function;
+	}
+    }
+
   /* Unparallelized OpenACC kernels constructs must get launched as 1 x 1 x 1
      kernels, so remove the parallelism dimensions function attributes
      potentially set earlier on.  */
diff --git a/gcc/testsuite/c-c++-common/goacc/classify-routine-nohost.c b/gcc/testsuite/c-c++-common/goacc/classify-routine-nohost.c
new file mode 100644
index 00000000000..a58482f7f92
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/classify-routine-nohost.c
@@ -0,0 +1,41 @@
+/* Check offloaded function's attributes and classification for OpenACC
+   routine with 'nohost' clause.  */
+
+/* { dg-additional-options "-O2" }
+   { dg-additional-options "-fopt-info-optimized-omp" }
+   { dg-additional-options "-fdump-tree-ompexp" }
+   { dg-additional-options "-fdump-tree-oaccdevlow" } */
+
+/* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
+   aspects of that functionality.  */
+
+#define N 1024
+
+extern unsigned int *__restrict a;
+extern unsigned int *__restrict b;
+extern unsigned int *__restrict c;
+#pragma acc declare copyin (a, b) create (c)
+
+#pragma acc routine nohost worker
+void ROUTINE ()
+{
+#pragma acc loop /* { dg-bogus "assigned OpenACC .* loop parallelism" } */
+  for (unsigned int i = 0; i < N; i++)
+    c[i] = a[i] + b[i];
+}
+
+/* Check the offloaded function's attributes.
+   { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(omp declare target \\(nohost worker\\), oacc function \\(0 1, 1 0, 1 0\\)\\)\\)" 1 "ompexp" } } */
+
+/* Check the offloaded function's classification.
+   { dg-final { scan-tree-dump-times "(?n)Function is OpenACC routine level 1" 1 "oaccdevlow" } }
+   { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'ROUTINE' has 'nohost' clause" 1 "oaccdevlow" { target c } } }
+   { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'void ROUTINE\\(\\)' has 'nohost' clause" 1 "oaccdevlow" { target { c++ && { ! offloading_enabled } } } } }
+   { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'ROUTINE\\(\\)' has 'nohost' clause" 1 "oaccdevlow" { target { c++ && offloading_enabled } } } }
+   { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'ROUTINE' discarded" 1 "oaccdevlow" { target c } } }
+   { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'void ROUTINE\\(\\)' discarded" 1 "oaccdevlow" { target { c++ && { ! offloading_enabled } } } } }
+   { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'ROUTINE\\(\\)' discarded" 1 "oaccdevlow" { target { c++ && offloading_enabled } } } }
+   TODO See PR101551 for 'offloading_enabled' differences.
+   { dg-final { scan-tree-dump-not "(?n)Compute dimensions" "oaccdevlow" } }
+   { dg-final { scan-tree-dump-not "(?n)__attribute__\\(.*omp declare target \\(nohost" "oaccdevlow" } }
+   { dg-final { scan-tree-dump-not "(?n)void ROUTINE \\(\\)" "oaccdevlow" } } */
diff --git a/gcc/testsuite/c-c++-common/goacc/classify-routine.c b/gcc/testsuite/c-c++-common/goacc/classify-routine.c
index 81fe3696baa..cc0ba2b9a7d 100644
--- a/gcc/testsuite/c-c++-common/goacc/classify-routine.c
+++ b/gcc/testsuite/c-c++-common/goacc/classify-routine.c
@@ -30,5 +30,13 @@ void ROUTINE ()
 /* Check the offloaded function's classification and compute dimensions (will
    always be 1 x 1 x 1 for non-offloading compilation).
    { dg-final { scan-tree-dump-times "(?n)Function is OpenACC routine level 1" 1 "oaccdevlow" } }
+   { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'ROUTINE' doesn't have 'nohost' clause" 1 "oaccdevlow" { target c } } }
+   { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'void ROUTINE\\(\\)' doesn't have 'nohost' clause" 1 "oaccdevlow" { target { c++ && { ! offloading_enabled } } } } }
+   { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'ROUTINE\\(\\)' doesn't have 'nohost' clause" 1 "oaccdevlow" { target { c++ && offloading_enabled } } } }
+   { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'ROUTINE' not discarded" 1 "oaccdevlow" { target c } } }
+   { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'void ROUTINE\\(\\)' not discarded" 1 "oaccdevlow" { target { c++ && { ! offloading_enabled } } } } }
+   { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'ROUTINE\\(\\)' not discarded" 1 "oaccdevlow" { target { c++ && offloading_enabled } } } }
+   TODO See PR101551 for 'offloading_enabled' differences.
    { dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccdevlow" } }
-   { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(0 1, 1 1, 1 1\\), omp declare target \\(worker\\), oacc function \\(0 1, 1 0, 1 0\\)\\)\\)" 1 "oaccdevlow" } } */
+   { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(0 1, 1 1, 1 1\\), omp declare target \\(worker\\), oacc function \\(0 1, 1 0, 1 0\\)\\)\\)" 1 "oaccdevlow" } }
+   { dg-final { scan-tree-dump-times "(?n)void ROUTINE \\(\\)" 1 "oaccdevlow" } } */
diff --git a/gcc/testsuite/c-c++-common/goacc/routine-2.c b/gcc/testsuite/c-c++-common/goacc/routine-2.c
index be1510a369c..3bf33e83d56 100644
--- a/gcc/testsuite/c-c++-common/goacc/routine-2.c
+++ b/gcc/testsuite/c-c++-common/goacc/routine-2.c
@@ -1,3 +1,7 @@
 /* Test invalid use of the OpenACC 'routine' directive.  */
 
 #pragma acc routine (nothing) gang /* { dg-error "not been declared" } */
+
+
+#pragma acc routine nohost nohost /* { dg-error "too many 'nohost' clauses" } */
+extern void nohost (void);
diff --git a/gcc/testsuite/c-c++-common/goacc/routine-nohost-1.c b/gcc/testsuite/c-c++-common/goacc/routine-nohost-1.c
new file mode 100644
index 00000000000..c8927416efa
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/routine-nohost-1.c
@@ -0,0 +1,50 @@
+/* Test OpenACC 'routine' with 'nohost' clause, valid use.  */
+
+/* { dg-additional-options "-fdump-tree-oaccdevlow" } */
+
+#pragma acc routine nohost
+int THREE(void)
+{
+  return 3;
+}
+
+#pragma acc routine (THREE) nohost
+
+#pragma acc routine nohost
+extern int THREE(void);
+
+/* { dg-final { scan-tree-dump-times {(?n)^OpenACC routine '[^']*THREE[^']*' has 'nohost' clause\.$} 1 oaccdevlow } } */
+
+
+#pragma acc routine nohost
+extern void NOTHING(void);
+
+#pragma acc routine (NOTHING) nohost
+
+void NOTHING(void)
+{
+}
+
+#pragma acc routine nohost
+extern void NOTHING(void);
+
+#pragma acc routine (NOTHING) nohost
+
+/* { dg-final { scan-tree-dump-times {(?n)^OpenACC routine '[^']*NOTHING[^']*' has 'nohost' clause\.$} 1 oaccdevlow } } */
+
+
+extern float ADD(float, float);
+
+#pragma acc routine (ADD) nohost
+
+float ADD(float x, float y)
+{
+  return x + y;
+}
+
+#pragma acc routine nohost
+extern float ADD(float, float);
+
+#pragma acc routine (ADD) nohost
+
+/* { dg-final { scan-tree-dump-times {(?n)^OpenACC routine '[^']*ADD[^']*' has 'nohost' clause\.$} 1 oaccdevlow } } */
diff --git a/gcc/testsuite/c-c++-common/goacc/routine-nohost-2.c b/gcc/testsuite/c-c++-common/goacc/routine-nohost-2.c
new file mode 100644
index 00000000000..d9acb805d2d
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/routine-nohost-2.c
@@ -0,0 +1,96 @@
+/* Test OpenACC 'routine' with 'nohost' clause, invalid use.  */
+
+#pragma acc routine /* { dg-note {\.\.\. without 'nohost' clause near to here} } */
+int THREE_1(void)
+{
+  return 3;
+}
+
+#pragma acc routine (THREE_1) \
+  nohost /* { dg-error {incompatible 'nohost' clause when applying '#pragma acc routine' to '[^']*THREE_1[^']*', which has already been marked with an OpenACC 'routine' directive} } */
+
+#pragma acc routine \
+  nohost /* { dg-error {incompatible 'nohost' clause when applying '#pragma acc routine' to '[^']*THREE_1[^']*', which has already been marked with an OpenACC 'routine' directive} } */
+extern int THREE_1(void);
+
+
+#pragma acc routine /* { dg-note {\.\.\. without 'nohost' clause near to here} } */
+extern void NOTHING_1(void);
+
+#pragma acc routine (NOTHING_1) \
+  nohost /* { dg-error {incompatible 'nohost' clause when applying '#pragma acc routine' to '[^']*NOTHING_1[^']*', which has already been marked with an OpenACC 'routine' directive} } */
+
+void NOTHING_1(void)
+{
+}
+
+#pragma acc routine \
+  nohost /* { dg-error {incompatible 'nohost' clause when applying '#pragma acc routine' to '[^']*NOTHING_1[^']*', which has already been marked with an OpenACC 'routine' directive} } */
+extern void NOTHING_1(void);
+
+#pragma acc routine (NOTHING_1) \
+  nohost /* { dg-error {incompatible 'nohost' clause when applying '#pragma acc routine' to '[^']*NOTHING_1[^']*', which has already been marked with an OpenACC 'routine' directive} } */
+
+
+extern float ADD_1(float, float);
+
+#pragma acc routine (ADD_1) /* { dg-note {\.\.\. without 'nohost' clause near to here} } */
+
+float ADD_1(float x, float y)
+{
+  return x + y;
+}
+
+#pragma acc routine \
+  nohost /* { dg-error {incompatible 'nohost' clause when applying '#pragma acc routine' to '[^']*ADD_1[^']*', which has already been marked with an OpenACC 'routine' directive} } */
+extern float ADD_1(float, float);
+
+#pragma acc routine (ADD_1) \
+  nohost /* { dg-error {incompatible 'nohost' clause when applying '#pragma acc routine' to '[^']*ADD_1[^']*', which has already been marked with an OpenACC 'routine' directive} } */
+
+
+/* The same again, but with/without nohost reversed.  */
+
+#pragma acc routine \
+  nohost /* { dg-note {\.\.\. with 'nohost' clause here} } */
+int THREE_2(void)
+{
+  return 3;
+}
+
+#pragma acc routine (THREE_2) /* { dg-error {missing 'nohost' clause when applying '#pragma acc routine' to '[^']*THREE_2[^']*', which has already been marked with an OpenACC 'routine' directive} } */
+
+#pragma acc routine /* { dg-error {missing 'nohost' clause when applying '#pragma acc routine' to '[^']*THREE_2[^']*', which has already been marked with an OpenACC 'routine' directive} } */
+extern int THREE_2(void);
+
+
+#pragma acc routine \
+  nohost /* { dg-note {\.\.\. with 'nohost' clause here} } */
+extern void NOTHING_2(void);
+
+#pragma acc routine (NOTHING_2) /* { dg-error {missing 'nohost' clause when applying '#pragma acc routine' to '[^']*NOTHING_2[^']*', which has already been marked with an OpenACC 'routine' directive} } */
+
+void NOTHING_2(void)
+{
+}
+
+#pragma acc routine /* { dg-error {missing 'nohost' clause when applying '#pragma acc routine' to '[^']*NOTHING_2[^']*', which has already been marked with an OpenACC 'routine' directive} } */
+extern void NOTHING_2(void);
+
+#pragma acc routine (NOTHING_2) /* { dg-error {missing 'nohost' clause when applying '#pragma acc routine' to '[^']*NOTHING_2[^']*', which has already been marked with an OpenACC 'routine' directive} } */
+
+
+extern float ADD_2(float, float);
+
+#pragma acc routine (ADD_2) \
+  nohost /* { dg-note {\.\.\. with 'nohost' clause here} } */
+
+float ADD_2(float x, float y)
+{
+  return x + y;
+}
+
+#pragma acc routine /* { dg-error {missing 'nohost' clause when applying '#pragma acc routine' to '[^']*ADD_2[^']*', which has already been marked with an OpenACC 'routine' directive} } */
+extern float ADD_2(float, float);
+
+#pragma acc routine (ADD_2) /* { dg-error {missing 'nohost' clause when applying '#pragma acc routine' to '[^']*ADD_2[^']*', which has already been marked with an OpenACC 'routine' directive} } */
diff --git a/gcc/testsuite/g++.dg/goacc/template.C b/gcc/testsuite/g++.dg/goacc/template.C
index 51a3f54e43f..f34fcfea52d 100644
--- a/gcc/testsuite/g++.dg/goacc/template.C
+++ b/gcc/testsuite/g++.dg/goacc/template.C
@@ -1,4 +1,6 @@
-#pragma acc routine
+/* { dg-additional-options "-fdump-tree-oaccdevlow" } */
+
+#pragma acc routine nohost
 template <typename T> T
 accDouble(int val)
 {
@@ -153,3 +155,14 @@ main ()
 
   return b + c;
 }
+
+/* { dg-final { scan-tree-dump-times {(?n)^OpenACC routine '[^']+' has 'nohost' clause\.$} 4 oaccdevlow } }
+   { dg-final { scan-tree-dump-times {(?n)^OpenACC routine 'T accDouble\(int\) \[with T = char\]' has 'nohost' clause\.$} 1 oaccdevlow { target { ! offloading_enabled } } } }
+   { dg-final { scan-tree-dump-times {(?n)^OpenACC routine 'accDouble<char>\(int\)char' has 'nohost' clause\.$} 1 oaccdevlow { target offloading_enabled } } }
+   { dg-final { scan-tree-dump-times {(?n)^OpenACC routine 'T accDouble\(int\) \[with T = int\]' has 'nohost' clause\.$} 1 oaccdevlow { target { ! offloading_enabled } } } }
+   { dg-final { scan-tree-dump-times {(?n)^OpenACC routine 'accDouble<int>\(int\)int' has 'nohost' clause\.$} 1 oaccdevlow { target offloading_enabled } } }
+   { dg-final { scan-tree-dump-times {(?n)^OpenACC routine 'T accDouble\(int\) \[with T = float\]' has 'nohost' clause\.$} 1 oaccdevlow { target { ! offloading_enabled } } } }
+   { dg-final { scan-tree-dump-times {(?n)^OpenACC routine 'accDouble<float>\(int\)float' has 'nohost' clause\.$} 1 oaccdevlow { target offloading_enabled } } }
+   { dg-final { scan-tree-dump-times {(?n)^OpenACC routine 'T accDouble\(int\) \[with T = double\]' has 'nohost' clause\.$} 1 oaccdevlow { target { ! offloading_enabled } } } }
+   { dg-final { scan-tree-dump-times {(?n)^OpenACC routine 'accDouble<double>\(int\)double' has 'nohost' clause\.$} 1 oaccdevlow { target offloading_enabled } } }
+   TODO See PR101551 for 'offloading_enabled' differences.  */
diff --git a/gcc/testsuite/gfortran.dg/goacc/classify-routine-nohost.f95 b/gcc/testsuite/gfortran.dg/goacc/classify-routine-nohost.f95
new file mode 100644
index 00000000000..0e06fb9f0ba
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/classify-routine-nohost.f95
@@ -0,0 +1,39 @@
+! Check offloaded function's attributes and classification for OpenACC
+! routine with 'nohost' clause.
+
+! { dg-additional-options "-O2" }
+! { dg-additional-options "-fopt-info-optimized-omp" }
+! { dg-additional-options "-fdump-tree-ompexp" }
+! { dg-additional-options "-fdump-tree-oaccdevlow" }
+
+! { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
+! aspects of that functionality.
+
+subroutine ROUTINE
+  !$acc routine nohost worker
+  integer, parameter :: n = 1024
+  integer, dimension (0:n-1) :: a, b, c
+  integer :: i
+
+  call setup(a, b)
+
+  !$acc loop ! { dg-bogus "assigned OpenACC .* loop parallelism" }
+  do i = 0, n - 1
+     c(i) = a(i) + b(i)
+  end do
+end subroutine ROUTINE
+
+! Check the offloaded function's attributes.
+! { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(0 1, 1 0, 1 0\\), omp declare target \\(nohost worker\\)\\)\\)" 1 "ompexp" } }
+
+! Check the offloaded function's classification.
+! { dg-final { scan-tree-dump-times "(?n)Function is OpenACC routine level 1" 1 "oaccdevlow" } }
+! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'routine' has 'nohost' clause" 1 "oaccdevlow" { target { ! offloading_enabled } } } }
+! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'routine_' has 'nohost' clause" 1 "oaccdevlow" { target offloading_enabled } } }
+! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'routine' discarded" 1 "oaccdevlow" { target { ! offloading_enabled } } } }
+! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'routine_' discarded" 1 "oaccdevlow" { target offloading_enabled } } }
+! { dg-final { scan-tree-dump-not "(?n)Compute dimensions" "oaccdevlow" } }
+! { dg-final { scan-tree-dump-not "(?n)__attribute__\\(.*omp declare target \\(nohost" "oaccdevlow" } }
+! { dg-final { scan-tree-dump-not "(?n)void routine \\(\\)" "oaccdevlow" { target { ! offloading_enabled } } } }
+! { dg-final { scan-tree-dump-not "(?n)void routine_ \\(\\)" "oaccdevlow" { target offloading_enabled } } }
+!TODO See PR101551 for 'offloading_enabled' differences.
diff --git a/gcc/testsuite/gfortran.dg/goacc/classify-routine.f95 b/gcc/testsuite/gfortran.dg/goacc/classify-routine.f95
index 52cc870dfba..92d3243cdcf 100644
--- a/gcc/testsuite/gfortran.dg/goacc/classify-routine.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/classify-routine.f95
@@ -29,5 +29,12 @@ end subroutine ROUTINE
 ! Check the offloaded function's classification and compute dimensions (will
 ! always be 1 x 1 x 1 for non-offloading compilation).
 ! { dg-final { scan-tree-dump-times "(?n)Function is OpenACC routine level 1" 1 "oaccdevlow" } }
+! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'routine' doesn't have 'nohost' clause" 1 "oaccdevlow" { target { ! offloading_enabled } } } }
+! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'routine_' doesn't have 'nohost' clause" 1 "oaccdevlow" { target offloading_enabled } } }
+! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'routine' not discarded" 1 "oaccdevlow" { target { ! offloading_enabled } } } }
+! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'routine_' not discarded" 1 "oaccdevlow" { target offloading_enabled } } }
 ! { dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccdevlow" } }
 ! { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(0 1, 1 1, 1 1\\), omp declare target \\(worker\\)\\)\\)" 1 "oaccdevlow" } }
+! { dg-final { scan-tree-dump-times "(?n)void routine \\(\\)" 1 "oaccdevlow" { target { ! offloading_enabled } } } }
+! { dg-final { scan-tree-dump-times "(?n)void routine_ \\(\\)" 1 "oaccdevlow" { target offloading_enabled } } }
+!TODO See PR101551 for 'offloading_enabled' differences.
diff --git a/gcc/testsuite/gfortran.dg/goacc/pure-elemental-procedures-2.f90 b/gcc/testsuite/gfortran.dg/goacc/pure-elemental-procedures-2.f90
index 97d92c3becc..31233b35fa7 100644
--- a/gcc/testsuite/gfortran.dg/goacc/pure-elemental-procedures-2.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/pure-elemental-procedures-2.f90
@@ -2,6 +2,10 @@ pure elemental subroutine foo()
 !$acc routine vector  ! { dg-error "ROUTINE with GANG, WORKER, or VECTOR clause is not permitted in PURE procedure" }
 end
 
+pure elemental subroutine foo_nh()
+!$acc routine nohost vector  ! { dg-error "ROUTINE with GANG, WORKER, or VECTOR clause is not permitted in PURE procedure" }
+end
+
 elemental subroutine foo2()
 !$acc routine (myfoo2) gang  ! { dg-error "Invalid NAME 'myfoo2' in" }
 end
@@ -10,18 +14,38 @@ elemental subroutine foo2a()
 !$acc routine gang  ! { dg-error "ROUTINE with GANG, WORKER, or VECTOR clause is not permitted in PURE procedure" }
 end
 
+elemental subroutine foo2a_nh()
+!$acc routine nohost gang  ! { dg-error "ROUTINE with GANG, WORKER, or VECTOR clause is not permitted in PURE procedure" }
+end
+
 pure subroutine foo3()
 !$acc routine vector ! { dg-error "ROUTINE with GANG, WORKER, or VECTOR clause is not permitted in PURE procedure" }
 end
 
+pure subroutine foo3_nh()
+!$acc routine nohost vector ! { dg-error "ROUTINE with GANG, WORKER, or VECTOR clause is not permitted in PURE procedure" }
+end
+
 elemental impure subroutine foo4()
 !$acc routine vector ! OK: impure
 end
 
+elemental impure subroutine foo4_nh()
+!$acc routine nohost vector ! OK: impure
+end
+
 pure subroutine foo5()
 !$acc routine seq ! OK: seq
 end
 
+pure subroutine foo5_nh()
+!$acc routine nohost seq ! OK: seq
+end
+
 pure subroutine foo6()
 !$acc routine ! OK (implied 'seq')
 end
+
+pure subroutine foo6_nh()
+!$acc routine nohost ! OK (implied 'seq')
+end
diff --git a/gcc/testsuite/gfortran.dg/goacc/routine-6.f90 b/gcc/testsuite/gfortran.dg/goacc/routine-6.f90
index f1e2aa3c3c3..3cd543e5aad 100644
--- a/gcc/testsuite/gfortran.dg/goacc/routine-6.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/routine-6.f90
@@ -116,3 +116,13 @@ subroutine subr10 (x)
      x = x * x - 1
   end if
 end subroutine subr10
+
+subroutine subr20 (x)
+  !$acc routine (subr20) nohost nohost ! { dg-error "Failed to match clause" }
+  integer, intent(inout) :: x
+  if (x < 1) then
+     x = 1
+  else
+     x = x * x - 1
+  end if
+end subroutine subr20
diff --git a/gcc/testsuite/gfortran.dg/goacc/routine-intrinsic-2.f b/gcc/testsuite/gfortran.dg/goacc/routine-intrinsic-2.f
index 22524cc1645..0372e7839e6 100644
--- a/gcc/testsuite/gfortran.dg/goacc/routine-intrinsic-2.f
+++ b/gcc/testsuite/gfortran.dg/goacc/routine-intrinsic-2.f
@@ -7,6 +7,11 @@
 !$ACC ROUTINE (ABORT) GANG ! { dg-error "Intrinsic symbol specified in \\!\\\$ACC ROUTINE \\( NAME \\) at \\(1\\) marked with incompatible GANG, WORKER, or VECTOR clause" }
 !$ACC ROUTINE (ABORT) VECTOR ! { dg-error "Intrinsic symbol specified in \\!\\\$ACC ROUTINE \\( NAME \\) at \\(1\\) marked with incompatible GANG, WORKER, or VECTOR clause" }
 
+!$ACC ROUTINE (ABORT) NOHOST ! { dg-error "Intrinsic symbol specified in \\!\\\$ACC ROUTINE \\( NAME \\) at \\(1\\) marked with incompatible NOHOST clause" }
+
+!$ACC ROUTINE (ABORT) WORKER NOHOST ! { dg-error "Intrinsic symbol specified in \\!\\\$ACC ROUTINE \\( NAME \\) at \\(1\\) marked with incompatible GANG, WORKER, or VECTOR clause" }
+!$ACC ROUTINE (ABORT) NOHOST GANG ! { dg-error "Intrinsic symbol specified in \\!\\\$ACC ROUTINE \\( NAME \\) at \\(1\\) marked with incompatible GANG, WORKER, or VECTOR clause" }
+
       CALL ABORT
       END SUBROUTINE sub_1
 
@@ -16,6 +21,11 @@
 !$ACC ROUTINE (ABORT) WORKER ! { dg-error "Intrinsic symbol specified in \\!\\\$ACC ROUTINE \\( NAME \\) at \\(1\\) marked with incompatible GANG, WORKER, or VECTOR clause" }
 !$ACC ROUTINE (ABORT) GANG ! { dg-error "Intrinsic symbol specified in \\!\\\$ACC ROUTINE \\( NAME \\) at \\(1\\) marked with incompatible GANG, WORKER, or VECTOR clause" }
 
+!$ACC ROUTINE (ABORT) NOHOST ! { dg-error "Intrinsic symbol specified in \\!\\\$ACC ROUTINE \\( NAME \\) at \\(1\\) marked with incompatible NOHOST clause" }
+
+!$ACC ROUTINE (ABORT) VECTOR NOHOST ! { dg-error "Intrinsic symbol specified in \\!\\\$ACC ROUTINE \\( NAME \\) at \\(1\\) marked with incompatible GANG, WORKER, or VECTOR clause" }
+!$ACC ROUTINE (ABORT) NOHOST WORKER ! { dg-error "Intrinsic symbol specified in \\!\\\$ACC ROUTINE \\( NAME \\) at \\(1\\) marked with incompatible GANG, WORKER, or VECTOR clause" }
+
       CONTAINS
       SUBROUTINE sub_2
       CALL ABORT
diff --git a/gcc/testsuite/gfortran.dg/goacc/routine-module-1.f90 b/gcc/testsuite/gfortran.dg/goacc/routine-module-1.f90
index 4e81f11fec8..46eec3d7488 100644
--- a/gcc/testsuite/gfortran.dg/goacc/routine-module-1.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/routine-module-1.f90
@@ -14,34 +14,48 @@ program main
   !$acc parallel loop seq ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
   do i = 1, 10
      call s_1 ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+     call s_1_nh ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
      call s_2 ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+     call s_2_nh ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
      call g_1 ! { dg-message "optimized: assigned OpenACC gang worker vector loop parallelism" }
+     call g_1_nh ! { dg-message "optimized: assigned OpenACC gang worker vector loop parallelism" }
      call w_1 ! { dg-message "optimized: assigned OpenACC worker vector loop parallelism" }
+     call w_1_nh ! { dg-message "optimized: assigned OpenACC worker vector loop parallelism" }
      call v_1 ! { dg-message "optimized: assigned OpenACC vector loop parallelism" }
+     call v_1_nh ! { dg-message "optimized: assigned OpenACC vector loop parallelism" }
   end do
   !$acc end parallel loop
 
   !$acc parallel loop gang ! { dg-message "optimized: assigned OpenACC gang loop parallelism" }
   do i = 1, 10
      call s_1 ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+     call s_1_nh ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
      call s_2 ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+     call s_2_nh ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
      call w_1 ! { dg-message "optimized: assigned OpenACC worker vector loop parallelism" }
+     call w_1_nh ! { dg-message "optimized: assigned OpenACC worker vector loop parallelism" }
      call v_1 ! { dg-message "optimized: assigned OpenACC vector loop parallelism" }
+     call v_1_nh ! { dg-message "optimized: assigned OpenACC vector loop parallelism" }
   end do
   !$acc end parallel loop
 
   !$acc parallel loop worker ! { dg-message "optimized: assigned OpenACC worker loop parallelism" }
   do i = 1, 10
      call s_1 ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+     call s_1_nh ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
      call s_2 ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+     call s_2_nh ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
      call v_1 ! { dg-message "optimized: assigned OpenACC vector loop parallelism" }
+     call v_1_nh ! { dg-message "optimized: assigned OpenACC vector loop parallelism" }
   end do
   !$acc end parallel loop
 
   !$acc parallel loop vector ! { dg-message "optimized: assigned OpenACC vector loop parallelism" }
   do i = 1, 10
      call s_1 ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+     call s_1_nh ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
      call s_2 ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+     call s_2_nh ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
   end do
   !$acc end parallel loop
 end program main
diff --git a/gcc/testsuite/gfortran.dg/goacc/routine-module-2.f90 b/gcc/testsuite/gfortran.dg/goacc/routine-module-2.f90
index eae0807643c..e796c1da300 100644
--- a/gcc/testsuite/gfortran.dg/goacc/routine-module-2.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/routine-module-2.f90
@@ -11,21 +11,27 @@ program main
   !$acc parallel loop gang
   do i = 1, 10
      call g_1 ! { dg-error "routine call uses same OpenACC parallelism as containing loop" }
+     call g_1_nh ! { dg-error "routine call uses same OpenACC parallelism as containing loop" }
   end do
   !$acc end parallel loop
 
   !$acc parallel loop worker
   do i = 1, 10
      call g_1 ! { dg-error "routine call uses same OpenACC parallelism as containing loop" }
+     call g_1_nh ! { dg-error "routine call uses same OpenACC parallelism as containing loop" }
      call w_1 ! { dg-error "routine call uses same OpenACC parallelism as containing loop" }
+     call w_1_nh ! { dg-error "routine call uses same OpenACC parallelism as containing loop" }
   end do
   !$acc end parallel loop
 
   !$acc parallel loop vector
   do i = 1, 10
      call g_1 ! { dg-error "routine call uses same OpenACC parallelism as containing loop" }
+     call g_1_nh ! { dg-error "routine call uses same OpenACC parallelism as containing loop" }
      call w_1 ! { dg-error "routine call uses same OpenACC parallelism as containing loop" }
+     call w_1_nh ! { dg-error "routine call uses same OpenACC parallelism as containing loop" }
      call v_1 ! { dg-error "routine call uses same OpenACC parallelism as containing loop" }
+     call v_1_nh ! { dg-error "routine call uses same OpenACC parallelism as containing loop" }
   end do
   !$acc end parallel loop
 end program main
diff --git a/gcc/testsuite/gfortran.dg/goacc/routine-module-3.f90 b/gcc/testsuite/gfortran.dg/goacc/routine-module-3.f90
index a4ff54954af..80fe07a3a91 100644
--- a/gcc/testsuite/gfortran.dg/goacc/routine-module-3.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/routine-module-3.f90
@@ -2,15 +2,54 @@
 
 ! { dg-compile-aux-modules "routine-module-mod-1.f90" }
 
-program main
+subroutine sr_1
   use routine_module_mod_1
   implicit none
+
   !$acc routine (s_1) seq ! { dg-error "Cannot change attributes of USE-associated symbol s_1" }
    ! { dg-error "NAME 's_1' invalid in \\!\\\$ACC ROUTINE \\( NAME \\)" "" { target *-*-* } .-1 }
+  !$acc routine (s_1_nh) seq nohost ! { dg-error "Cannot change attributes of USE-associated symbol s_1_nh" }
+   ! { dg-error "NAME 's_1_nh' invalid in \\!\\\$ACC ROUTINE \\( NAME \\)" "" { target *-*-* } .-1 }
   !$acc routine (s_2) seq ! { dg-error "Cannot change attributes of USE-associated symbol s_2" }
    ! { dg-error "NAME 's_2' invalid in \\!\\\$ACC ROUTINE \\( NAME \\)" "" { target *-*-* } .-1 }
+  !$acc routine (s_2_nh) seq nohost ! { dg-error "Cannot change attributes of USE-associated symbol s_2_nh" }
+   ! { dg-error "NAME 's_2_nh' invalid in \\!\\\$ACC ROUTINE \\( NAME \\)" "" { target *-*-* } .-1 }
   !$acc routine (v_1) seq ! { dg-error "Cannot change attributes of USE-associated symbol v_1" }
    ! { dg-error "NAME 'v_1' invalid in \\!\\\$ACC ROUTINE \\( NAME \\)" "" { target *-*-* } .-1 }
+  !$acc routine (v_1_nh) seq nohost ! { dg-error "Cannot change attributes of USE-associated symbol v_1_nh" }
+   ! { dg-error "NAME 'v_1_nh' invalid in \\!\\\$ACC ROUTINE \\( NAME \\)" "" { target *-*-* } .-1 }
   !$acc routine (w_1) gang ! { dg-error "Cannot change attributes of USE-associated symbol w_1" }
    ! { dg-error "NAME 'w_1' invalid in \\!\\\$ACC ROUTINE \\( NAME \\)" "" { target *-*-* } .-1 }
-end program main
+  !$acc routine (w_1_nh) gang nohost ! { dg-error "Cannot change attributes of USE-associated symbol w_1_nh" }
+   ! { dg-error "NAME 'w_1_nh' invalid in \\!\\\$ACC ROUTINE \\( NAME \\)" "" { target *-*-* } .-1 }
+  !$acc routine (g_1) gang ! { dg-error "Cannot change attributes of USE-associated symbol g_1" }
+   ! { dg-error "NAME 'g_1' invalid in \\!\\\$ACC ROUTINE \\( NAME \\)" "" { target *-*-* } .-1 }
+  !$acc routine (g_1_nh) gang nohost ! { dg-error "Cannot change attributes of USE-associated symbol g_1_nh" }
+   ! { dg-error "NAME 'g_1_nh' invalid in \\!\\\$ACC ROUTINE \\( NAME \\)" "" { target *-*-* } .-1 }
+end subroutine sr_1
+
+subroutine sr_2
+  use routine_module_mod_1
+  implicit none
+
+  !$acc routine (s_1) seq nohost ! { dg-error "Cannot change attributes of USE-associated symbol s_1" }
+   ! { dg-error "NAME 's_1' invalid in \\!\\\$ACC ROUTINE \\( NAME \\)" "" { target *-*-* } .-1 }
+  !$acc routine (s_1_nh) seq ! { dg-error "Cannot change attributes of USE-associated symbol s_1_nh" }
+   ! { dg-error "NAME 's_1_nh' invalid in \\!\\\$ACC ROUTINE \\( NAME \\)" "" { target *-*-* } .-1 }
+  !$acc routine (s_2) seq nohost ! { dg-error "Cannot change attributes of USE-associated symbol s_2" }
+   ! { dg-error "NAME 's_2' invalid in \\!\\\$ACC ROUTINE \\( NAME \\)" "" { target *-*-* } .-1 }
+  !$acc routine (s_2_nh) seq ! { dg-error "Cannot change attributes of USE-associated symbol s_2_nh" }
+   ! { dg-error "NAME 's_2_nh' invalid in \\!\\\$ACC ROUTINE \\( NAME \\)" "" { target *-*-* } .-1 }
+  !$acc routine (v_1) vector nohost ! { dg-error "Cannot change attributes of USE-associated symbol v_1" }
+   ! { dg-error "NAME 'v_1' invalid in \\!\\\$ACC ROUTINE \\( NAME \\)" "" { target *-*-* } .-1 }
+  !$acc routine (v_1_nh) vector ! { dg-error "Cannot change attributes of USE-associated symbol v_1_nh" }
+   ! { dg-error "NAME 'v_1_nh' invalid in \\!\\\$ACC ROUTINE \\( NAME \\)" "" { target *-*-* } .-1 }
+  !$acc routine (w_1) worker nohost ! { dg-error "Cannot change attributes of USE-associated symbol w_1" }
+   ! { dg-error "NAME 'w_1' invalid in \\!\\\$ACC ROUTINE \\( NAME \\)" "" { target *-*-* } .-1 }
+  !$acc routine (w_1_nh) worker ! { dg-error "Cannot change attributes of USE-associated symbol w_1_nh" }
+   ! { dg-error "NAME 'w_1_nh' invalid in \\!\\\$ACC ROUTINE \\( NAME \\)" "" { target *-*-* } .-1 }
+  !$acc routine (g_1) worker nohost ! { dg-error "Cannot change attributes of USE-associated symbol g_1" }
+   ! { dg-error "NAME 'g_1' invalid in \\!\\\$ACC ROUTINE \\( NAME \\)" "" { target *-*-* } .-1 }
+  !$acc routine (g_1_nh) worker ! { dg-error "Cannot change attributes of USE-associated symbol g_1_nh" }
+   ! { dg-error "NAME 'g_1_nh' invalid in \\!\\\$ACC ROUTINE \\( NAME \\)" "" { target *-*-* } .-1 }
+end subroutine sr_2
diff --git a/gcc/testsuite/gfortran.dg/goacc/routine-module-mod-1.f90 b/gcc/testsuite/gfortran.dg/goacc/routine-module-mod-1.f90
index 835619c6509..10e109675dc 100644
--- a/gcc/testsuite/gfortran.dg/goacc/routine-module-mod-1.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/routine-module-mod-1.f90
@@ -19,6 +19,17 @@ contains
     end do
   end subroutine s_1
 
+  subroutine s_1_nh
+    implicit none
+    !$acc routine nohost
+
+    integer :: i
+
+    !$acc loop ! { dg-bogus "assigned OpenACC .* loop parallelism" }
+    do i = 1, 3
+    end do
+  end subroutine s_1_nh
+
   subroutine s_2
     implicit none
     !$acc routine (s_2) seq
@@ -31,6 +42,17 @@ contains
     end do
   end subroutine s_2
 
+  subroutine s_2_nh
+    implicit none
+    !$acc routine (s_2_nh) seq nohost
+
+    integer :: i
+
+    !$acc loop ! { dg-bogus "assigned OpenACC .* loop parallelism" }
+    do i = 1, 3
+    end do
+  end subroutine s_2_nh
+
   subroutine v_1
     implicit none
     !$acc routine vector
@@ -42,6 +64,17 @@ contains
     end do
   end subroutine v_1
 
+  subroutine v_1_nh
+    implicit none
+    !$acc routine vector nohost
+
+    integer :: i
+
+    !$acc loop ! { dg-bogus "assigned OpenACC .* loop parallelism" }
+    do i = 1, 3
+    end do
+  end subroutine v_1_nh
+
   subroutine w_1
     implicit none
     !$acc routine (w_1) worker
@@ -53,6 +86,17 @@ contains
     end do
   end subroutine w_1
 
+  subroutine w_1_nh
+    implicit none
+    !$acc routine (w_1_nh) worker nohost
+
+    integer :: i
+
+    !$acc loop ! { dg-bogus "assigned OpenACC .* loop parallelism" }
+    do i = 1, 3
+    end do
+  end subroutine w_1_nh
+
   subroutine g_1
     implicit none
     !$acc routine gang
@@ -65,6 +109,17 @@ contains
     end do
   end subroutine g_1
 
+  subroutine g_1_nh
+    implicit none
+    !$acc routine gang nohost
+
+    integer :: i
+
+    !$acc loop ! { dg-bogus "assigned OpenACC .* loop parallelism" }
+    do i = 1, 3
+    end do
+  end subroutine g_1_nh
+
   subroutine pl_1
     implicit none
 
@@ -74,10 +129,15 @@ contains
     ! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
     do i = 1, 3
        call s_1 ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+       call s_1_nh ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
        call s_2 ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+       call s_2_nh ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
        call v_1 ! { dg-message "optimized: assigned OpenACC vector loop parallelism" }
+       call v_1_nh ! { dg-message "optimized: assigned OpenACC vector loop parallelism" }
        call w_1 ! { dg-message "optimized: assigned OpenACC worker vector loop parallelism" }
+       call w_1_nh ! { dg-message "optimized: assigned OpenACC worker vector loop parallelism" }
        call g_1 ! { dg-message "optimized: assigned OpenACC gang worker vector loop parallelism" }
+       call g_1_nh ! { dg-message "optimized: assigned OpenACC gang worker vector loop parallelism" }
     end do
   end subroutine pl_1
 end module routine_module_mod_1
diff --git a/gcc/testsuite/gfortran.dg/goacc/routine-multiple-directives-1.f90 b/gcc/testsuite/gfortran.dg/goacc/routine-multiple-directives-1.f90
index 622a9d9ccce..44ef4533f04 100644
--- a/gcc/testsuite/gfortran.dg/goacc/routine-multiple-directives-1.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/routine-multiple-directives-1.f90
@@ -1,5 +1,8 @@
 ! Check for valid cases of multiple OpenACC 'routine' directives.
 
+! { dg-additional-options "-fdump-tree-oaccdevlow" }
+!TODO See PR101551 for 'offloading_enabled' differences.
+
 ! { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
 ! aspects of that functionality.
 
@@ -8,12 +11,32 @@
 !$ACC ROUTINE(s_1) SEQ
 !$ACC ROUTINE SEQ
       END SUBROUTINE s_1
+      ! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 's_1' doesn't have 'nohost' clause" 1 "oaccdevlow" { target { ! offloading_enabled } } } }
+      ! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 's_1_' doesn't have 'nohost' clause" 1 "oaccdevlow" { target offloading_enabled } } }
+
+      SUBROUTINE s_1_nh
+!$ACC ROUTINE(s_1_nh) NOHOST
+!$ACC ROUTINE(s_1_nh) SEQ NOHOST
+!$ACC ROUTINE NOHOST SEQ
+      END SUBROUTINE s_1_nh
+      ! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 's_1_nh' has 'nohost' clause" 1 "oaccdevlow" { target { ! offloading_enabled } } } }
+      ! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 's_1_nh_' has 'nohost' clause" 1 "oaccdevlow" { target offloading_enabled } } }
 
       SUBROUTINE s_2
 !$ACC ROUTINE
 !$ACC ROUTINE SEQ
 !$ACC ROUTINE(s_2)
       END SUBROUTINE s_2
+      ! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 's_2' doesn't have 'nohost' clause" 1 "oaccdevlow" { target { ! offloading_enabled } } } }
+      ! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 's_2_' doesn't have 'nohost' clause" 1 "oaccdevlow" { target offloading_enabled } } }
+
+      SUBROUTINE s_2_nh
+!$ACC ROUTINE NOHOST
+!$ACC ROUTINE NOHOST SEQ
+!$ACC ROUTINE(s_2_nh) NOHOST
+      END SUBROUTINE s_2_nh
+      ! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 's_2_nh' has 'nohost' clause" 1 "oaccdevlow" { target { ! offloading_enabled } } } }
+      ! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 's_2_nh_' has 'nohost' clause" 1 "oaccdevlow" { target offloading_enabled } } }
 
       SUBROUTINE v_1
 !$ACC ROUTINE VECTOR
@@ -22,6 +45,18 @@
 !$ACC ROUTINE VECTOR
 ! { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-5 }
       END SUBROUTINE v_1
+      ! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'v_1' doesn't have 'nohost' clause" 1 "oaccdevlow" { target { ! offloading_enabled } } } }
+      ! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'v_1_' doesn't have 'nohost' clause" 1 "oaccdevlow" { target offloading_enabled } } }
+
+      SUBROUTINE v_1_nh
+!$ACC ROUTINE NOHOST VECTOR
+!$ACC ROUTINE VECTOR NOHOST
+!$ACC ROUTINE(v_1_nh) NOHOST VECTOR
+!$ACC ROUTINE VECTOR NOHOST
+! { dg-bogus "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-5 }
+      END SUBROUTINE v_1_nh
+      ! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'v_1_nh' has 'nohost' clause" 1 "oaccdevlow" { target { ! offloading_enabled } } } }
+      ! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'v_1_nh_' has 'nohost' clause" 1 "oaccdevlow" { target offloading_enabled } } }
 
       SUBROUTINE v_2
 !$ACC ROUTINE(v_2) VECTOR
@@ -29,6 +64,17 @@
 !$ACC ROUTINE(v_2) VECTOR
 ! { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-4 }
       END SUBROUTINE v_2
+      ! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'v_2' doesn't have 'nohost' clause" 1 "oaccdevlow" { target { ! offloading_enabled } } } }
+      ! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'v_2_' doesn't have 'nohost' clause" 1 "oaccdevlow" { target offloading_enabled } } }
+
+      SUBROUTINE v_2_nh
+!$ACC ROUTINE(v_2_nh) VECTOR NOHOST
+!$ACC ROUTINE VECTOR NOHOST
+!$ACC ROUTINE(v_2_nh) NOHOST VECTOR
+! { dg-bogus "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-4 }
+      END SUBROUTINE v_2_nh
+      ! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'v_2_nh' has 'nohost' clause" 1 "oaccdevlow" { target { ! offloading_enabled } } } }
+      ! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'v_2_nh_' has 'nohost' clause" 1 "oaccdevlow" { target offloading_enabled } } }
 
       SUBROUTINE sub_1
       IMPLICIT NONE
@@ -36,12 +82,22 @@
 !$ACC ROUTINE (g_1) GANG
 !$ACC ROUTINE (g_1) GANG
 !$ACC ROUTINE (g_1) GANG
+      EXTERNAL :: g_1_nh
+!$ACC ROUTINE (g_1_nh) GANG NOHOST
+!$ACC ROUTINE (g_1_nh) NOHOST GANG
+!$ACC ROUTINE (g_1_nh) NOHOST GANG
+!$ACC ROUTINE (g_1_nh) GANG NOHOST
 
       CALL s_1
+      CALL s_1_nh
       CALL s_2
+      CALL s_2_nh
       CALL v_1
+      CALL v_1_nh
       CALL v_2
+      CALL v_2_nh
       CALL g_1
+      CALL g_1_nh
       CALL ABORT
       END SUBROUTINE sub_1
 
@@ -50,14 +106,22 @@
       EXTERNAL :: w_1
 !$ACC ROUTINE (w_1) WORKER
 !$ACC ROUTINE (w_1) WORKER
+      EXTERNAL :: w_1_nh
+!$ACC ROUTINE (w_1_nh) NOHOST WORKER
+!$ACC ROUTINE (w_1_nh) WORKER NOHOST
 
       CONTAINS
       SUBROUTINE sub_2
       CALL s_1
+      CALL s_1_nh
       CALL s_2
+      CALL s_2_nh
       CALL v_1
+      CALL v_1_nh
       CALL v_2
+      CALL v_2_nh
       CALL w_1
+      CALL w_1_nh
       CALL ABORT
       END SUBROUTINE sub_2
       END MODULE m_w_1
diff --git a/gcc/testsuite/gfortran.dg/goacc/routine-multiple-directives-2.f90 b/gcc/testsuite/gfortran.dg/goacc/routine-multiple-directives-2.f90
index 54365ae3f4e..f332ed5bad3 100644
--- a/gcc/testsuite/gfortran.dg/goacc/routine-multiple-directives-2.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/routine-multiple-directives-2.f90
@@ -9,8 +9,32 @@
 !$ACC ROUTINE
 !$ACC ROUTINE(s_1) WORKER ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
 !$ACC ROUTINE GANG VECTOR ! { dg-error "Multiple loop axes specified for routine" }
+!$ACC ROUTINE VECTOR NOHOST WORKER ! { dg-error "Multiple loop axes specified for routine" }
+!$ACC ROUTINE(s_1) NOHOST ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE NOHOST GANG ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE(s_1) SEQ NOHOST ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE NOHOST ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE(s_1) NOHOST WORKER ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE GANG NOHOST VECTOR ! { dg-error "Multiple loop axes specified for routine" }
       END SUBROUTINE s_1
 
+      SUBROUTINE s_1_nh
+!$ACC ROUTINE NOHOST VECTOR WORKER ! { dg-error "Multiple loop axes specified for routine" }
+!$ACC ROUTINE(s_1_nh) NOHOST
+!$ACC ROUTINE NOHOST GANG ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE(s_1_nh) NOHOST SEQ
+!$ACC ROUTINE NOHOST
+!$ACC ROUTINE(s_1_nh) WORKER NOHOST ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE GANG NOHOST VECTOR ! { dg-error "Multiple loop axes specified for routine" }
+!$ACC ROUTINE VECTOR WORKER ! { dg-error "Multiple loop axes specified for routine" }
+!$ACC ROUTINE(s_1_nh) ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE GANG ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE(s_1_nh) SEQ ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE(s_1_nh) WORKER ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE GANG VECTOR ! { dg-error "Multiple loop axes specified for routine" }
+      END SUBROUTINE s_1_nh
+
       SUBROUTINE s_2
 !$ACC ROUTINE(s_2) VECTOR WORKER ! { dg-error "Multiple loop axes specified for routine" }
 !$ACC ROUTINE
@@ -19,8 +43,32 @@
 !$ACC ROUTINE(s_2)
 !$ACC ROUTINE WORKER ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
 !$ACC ROUTINE(s_2) GANG VECTOR ! { dg-error "Multiple loop axes specified for routine" }
+!$ACC ROUTINE(s_2) VECTOR NOHOST WORKER ! { dg-error "Multiple loop axes specified for routine" }
+!$ACC ROUTINE NOHOST ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE(s_2) GANG NOHOST ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE SEQ NOHOST ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE(s_2) NOHOST ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE NOHOST WORKER ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE(s_2) NOHOST GANG VECTOR ! { dg-error "Multiple loop axes specified for routine" }
       END SUBROUTINE s_2
 
+      SUBROUTINE s_2_nh
+!$ACC ROUTINE(s_2_nh) NOHOST VECTOR WORKER ! { dg-error "Multiple loop axes specified for routine" }
+!$ACC ROUTINE NOHOST
+!$ACC ROUTINE(s_2_nh) GANG NOHOST ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE SEQ NOHOST
+!$ACC ROUTINE(s_2_nh) NOHOST
+!$ACC ROUTINE NOHOST WORKER ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE(s_2_nh) NOHOST GANG VECTOR ! { dg-error "Multiple loop axes specified for routine" }
+!$ACC ROUTINE(s_2_nh) VECTOR WORKER ! { dg-error "Multiple loop axes specified for routine" }
+!$ACC ROUTINE ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE(s_2_nh) GANG ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE SEQ ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE(s_2_nh) ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE WORKER ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE(s_2_nh) GANG VECTOR ! { dg-error "Multiple loop axes specified for routine" }
+      END SUBROUTINE s_2_nh
+
       SUBROUTINE v_1
 !$ACC ROUTINE VECTOR WORKER ! { dg-error "Multiple loop axes specified for routine" }
 !$ACC ROUTINE VECTOR
@@ -30,16 +78,61 @@
 !$ACC ROUTINE(v_1) VECTOR
 !$ACC ROUTINE WORKER ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
 !$ACC ROUTINE GANG VECTOR ! { dg-error "Multiple loop axes specified for routine" }
+!$ACC ROUTINE NOHOST VECTOR WORKER ! { dg-error "Multiple loop axes specified for routine" }
+!$ACC ROUTINE NOHOST VECTOR ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE GANG NOHOST ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE NOHOST SEQ ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE NOHOST ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE(v_1) VECTOR NOHOST ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE WORKER NOHOST ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE GANG VECTOR NOHOST ! { dg-error "Multiple loop axes specified for routine" }
       END SUBROUTINE v_1
 
+      SUBROUTINE v_1_nh
+!$ACC ROUTINE VECTOR WORKER NOHOST ! { dg-error "Multiple loop axes specified for routine" }
+!$ACC ROUTINE VECTOR NOHOST
+!$ACC ROUTINE GANG NOHOST ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE NOHOST SEQ ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE NOHOST ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE(v_1_nh) VECTOR NOHOST
+!$ACC ROUTINE WORKER NOHOST ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE GANG NOHOST VECTOR ! { dg-error "Multiple loop axes specified for routine" }
+!$ACC ROUTINE VECTOR WORKER ! { dg-error "Multiple loop axes specified for routine" }
+!$ACC ROUTINE VECTOR ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE GANG ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE SEQ ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE(v_1_nh) VECTOR ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE WORKER ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE GANG VECTOR ! { dg-error "Multiple loop axes specified for routine" }
+      END SUBROUTINE v_1_nh
+
       SUBROUTINE v_2
 !$ACC ROUTINE(v_2) VECTOR
 !$ACC ROUTINE(v_2) VECTOR WORKER ! { dg-error "Multiple loop axes specified for routine" }
 !$ACC ROUTINE(v_2) ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
 !$ACC ROUTINE VECTOR
 !$ACC ROUTINE(v_2) GANG VECTOR ! { dg-error "Multiple loop axes specified for routine" }
+!$ACC ROUTINE(v_2) VECTOR NOHOST ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE(v_2) VECTOR NOHOST WORKER ! { dg-error "Multiple loop axes specified for routine" }
+!$ACC ROUTINE(v_2) NOHOST ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE VECTOR NOHOST ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE(v_2) NOHOST GANG VECTOR ! { dg-error "Multiple loop axes specified for routine" }
       END SUBROUTINE v_2
 
+      SUBROUTINE v_2_nh
+!$ACC ROUTINE(v_2_nh) VECTOR NOHOST
+!$ACC ROUTINE(v_2_nh) VECTOR WORKER NOHOST ! { dg-error "Multiple loop axes specified for routine" }
+!$ACC ROUTINE(v_2_nh) NOHOST ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE VECTOR NOHOST
+!$ACC ROUTINE(v_2_nh) GANG NOHOST VECTOR ! { dg-error "Multiple loop axes specified for routine" }
+!$ACC ROUTINE(v_2_nh) VECTOR ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE(v_2_nh) VECTOR WORKER ! { dg-error "Multiple loop axes specified for routine" }
+!$ACC ROUTINE(v_2_nh) ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE VECTOR ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE(v_2_nh) GANG VECTOR ! { dg-error "Multiple loop axes specified for routine" }
+      END SUBROUTINE v_2_nh
+
       SUBROUTINE sub_1
       IMPLICIT NONE
       EXTERNAL :: g_1
@@ -50,12 +143,39 @@
 !$ACC ROUTINE (g_1) ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
 !$ACC ROUTINE (g_1) GANG
 !$ACC ROUTINE (g_1) ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE (g_1) NOHOST GANG ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE (g_1) GANG WORKER NOHOST ! { dg-error "Multiple loop axes specified for routine" }
+!$ACC ROUTINE (g_1) NOHOST VECTOR ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE (g_1) NOHOST SEQ ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE (g_1) NOHOST ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE (g_1) GANG NOHOST ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE (g_1) NOHOST ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+      EXTERNAL :: g_1_nh
+!$ACC ROUTINE (g_1_nh) NOHOST GANG
+!$ACC ROUTINE (g_1_nh) GANG NOHOST WORKER ! { dg-error "Multiple loop axes specified for routine" }
+!$ACC ROUTINE (g_1_nh) NOHOST VECTOR ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE (g_1_nh) SEQ NOHOST ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE (g_1_nh) NOHOST ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE (g_1_nh) GANG NOHOST
+!$ACC ROUTINE (g_1_nh) NOHOST ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE (g_1_nh) GANG ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE (g_1_nh) GANG WORKER ! { dg-error "Multiple loop axes specified for routine" }
+!$ACC ROUTINE (g_1_nh) VECTOR ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE (g_1_nh) SEQ ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE (g_1_nh) ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE (g_1_nh) GANG ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE (g_1_nh) ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
 
       CALL s_1
+      CALL s_1_nh
       CALL s_2
+      CALL s_2_nh
       CALL v_1
+      CALL v_1_nh
       CALL v_2
+      CALL v_2_nh
       CALL g_1
+      CALL g_1_nh
       CALL ABORT
       END SUBROUTINE sub_1
 
@@ -69,14 +189,41 @@
 !$ACC ROUTINE (w_1) SEQ ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
 !$ACC ROUTINE (w_1) ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
 !$ACC ROUTINE (w_1) VECTOR ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE (w_1) WORKER NOHOST ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE (w_1) WORKER NOHOST SEQ ! { dg-error "Multiple loop axes specified for routine" }
+!$ACC ROUTINE (w_1) NOHOST ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE (w_1) NOHOST WORKER ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE (w_1) SEQ NOHOST ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE (w_1) NOHOST ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE (w_1) VECTOR NOHOST ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+      EXTERNAL :: w_1_nh
+!$ACC ROUTINE (w_1_nh) WORKER NOHOST
+!$ACC ROUTINE (w_1_nh) WORKER NOHOST SEQ ! { dg-error "Multiple loop axes specified for routine" }
+!$ACC ROUTINE (w_1_nh) NOHOST ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE (w_1_nh) NOHOST WORKER
+!$ACC ROUTINE (w_1_nh) NOHOST SEQ ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE (w_1_nh) NOHOST ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE (w_1_nh) VECTOR NOHOST ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE (w_1_nh) WORKER ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE (w_1_nh) WORKER SEQ ! { dg-error "Multiple loop axes specified for routine" }
+!$ACC ROUTINE (w_1_nh) ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE (w_1_nh) WORKER ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE (w_1_nh) SEQ ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE (w_1_nh) ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
+!$ACC ROUTINE (w_1_nh) VECTOR ! { dg-error "\\!\\\$ACC ROUTINE already applied" }
 
       CONTAINS
       SUBROUTINE sub_2
       CALL s_1
+      CALL s_1_nh
       CALL s_2
+      CALL s_2_nh
       CALL v_1
+      CALL v_1_nh
       CALL v_2
+      CALL v_2_nh
       CALL w_1
+      CALL w_1_nh
       CALL ABORT
       END SUBROUTINE sub_2
       END MODULE m_w_1
diff --git a/gcc/tree-core.h b/gcc/tree-core.h
index 93916090432..bfab988ecdd 100644
--- a/gcc/tree-core.h
+++ b/gcc/tree-core.h
@@ -508,7 +508,10 @@ enum omp_clause_code {
   OMP_CLAUSE_IF_PRESENT,
 
   /* OpenACC clause: finalize.  */
-  OMP_CLAUSE_FINALIZE
+  OMP_CLAUSE_FINALIZE,
+
+  /* OpenACC clause: nohost.  */
+  OMP_CLAUSE_NOHOST,
 };
 
 #undef DEFTREESTRUCT
diff --git a/gcc/tree-nested.c b/gcc/tree-nested.c
index 9edd922a303..0c3fb029054 100644
--- a/gcc/tree-nested.c
+++ b/gcc/tree-nested.c
@@ -1510,6 +1510,9 @@ convert_nonlocal_omp_clauses (tree *pclauses, struct walk_stmt_info *wi)
 	case OMP_CLAUSE__REDUCTEMP_:
 	case OMP_CLAUSE__SIMDUID_:
 	case OMP_CLAUSE__SIMT_:
+	  /* The following clauses are only allowed on OpenACC 'routine'
+	     directives, not seen here.  */
+	case OMP_CLAUSE_NOHOST:
 	  /* Anything else.  */
 	default:
 	  gcc_unreachable ();
@@ -2291,6 +2294,9 @@ convert_local_omp_clauses (tree *pclauses, struct walk_stmt_info *wi)
 	case OMP_CLAUSE__REDUCTEMP_:
 	case OMP_CLAUSE__SIMDUID_:
 	case OMP_CLAUSE__SIMT_:
+	  /* The following clauses are only allowed on OpenACC 'routine'
+	     directives, not seen here.  */
+	case OMP_CLAUSE_NOHOST:
 	  /* Anything else.  */
 	default:
 	  gcc_unreachable ();
diff --git a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c
index fde07dfd0e1..7201bd7d9f6 100644
--- a/gcc/tree-pretty-print.c
+++ b/gcc/tree-pretty-print.c
@@ -1303,6 +1303,9 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
     case OMP_CLAUSE_FINALIZE:
       pp_string (pp, "finalize");
       break;
+    case OMP_CLAUSE_NOHOST:
+      pp_string (pp, "nohost");
+      break;
     case OMP_CLAUSE_DETACH:
       pp_string (pp, "detach(");
       dump_generic_node (pp, OMP_CLAUSE_DECL (clause), spc, flags,
diff --git a/gcc/tree.c b/gcc/tree.c
index bead1ac134c..c621f870880 100644
--- a/gcc/tree.c
+++ b/gcc/tree.c
@@ -361,6 +361,7 @@ unsigned const char omp_clause_num_ops[] =
   3, /* OMP_CLAUSE_TILE  */
   0, /* OMP_CLAUSE_IF_PRESENT */
   0, /* OMP_CLAUSE_FINALIZE */
+  0, /* OMP_CLAUSE_NOHOST */
 };
 
 const char * const omp_clause_code_name[] =
@@ -448,6 +449,7 @@ const char * const omp_clause_code_name[] =
   "tile",
   "if_present",
   "finalize",
+  "nohost",
 };
 
 
@@ -11165,6 +11167,7 @@ walk_tree_1 (tree *tp, walk_tree_fn func, void *data,
 	case OMP_CLAUSE__SIMT_:
 	case OMP_CLAUSE_IF_PRESENT:
 	case OMP_CLAUSE_FINALIZE:
+	case OMP_CLAUSE_NOHOST:
 	  WALK_SUBTREE_TAIL (OMP_CLAUSE_CHAIN (*tp));
 
 	case OMP_CLAUSE_LASTPRIVATE:
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-nohost-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-nohost-1.c
new file mode 100644
index 00000000000..dc92727d5be
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-nohost-1.c
@@ -0,0 +1,63 @@
+/* Test 'nohost' clause via 'acc_on_device'.
+
+   With optimizations disabled, we currently don't expect that 'acc_on_device' "evaluates at compile time to a constant".
+   { dg-skip-if "TODO PR82391" { *-*-* } { "-O0" } }
+*/
+
+/* { dg-additional-options "-fdump-tree-oaccdevlow" } */
+
+/* { dg-additional-options "-fno-inline" } for stable results regarding OpenACC 'routine'.  */
+
+#include <assert.h>
+#include <openacc.h>
+
+#pragma acc routine
+static int fact(int n)
+{
+  if (n == 0 || n == 1)
+    return 1;
+  else
+    return n * fact(n - 1);
+}
+
+#pragma acc routine nohost
+static int fact_nohost(int n)
+{
+  /* Make sure this fails host compilation.  */
+#if defined ACC_DEVICE_TYPE_host
+  asm ("IT'S A TRAP");
+#elif defined ACC_DEVICE_TYPE_nvidia
+  asm ("{\n\t  .reg .u32 %tid_x;\n\t  mov.u32 %tid_x, %tid.x;\n\t}");
+#elif defined ACC_DEVICE_TYPE_radeon
+  asm ("s_nop 0");
+#else
+# error Not ported to this ACC_DEVICE_TYPE
+#endif
+
+  return fact(n);
+}
+/* { dg-final { scan-tree-dump-times {(?n)^OpenACC routine 'fact_nohost' has 'nohost' clause\.$} 1 oaccdevlow { target c } } }
+   { dg-final { scan-tree-dump-times {(?n)^OpenACC routine 'int fact_nohost\(int\)' has 'nohost' clause\.$} 1 oaccdevlow { target { c++ && { ! offloading_enabled } } } } }
+   { dg-final { scan-tree-dump-times {(?n)^OpenACC routine 'fact_nohost\(int\)' has 'nohost' clause\.$} 1 oaccdevlow { target { c++ && offloading_enabled } } } }
+   TODO See PR101551 for 'offloading_enabled' differences.  */
+
+int main()
+{
+#define N 10
+  int x[N];
+
+#pragma acc parallel loop copyout(x)
+  for (int i = 0; i < N; ++i)
+    /*TODO PR82391: '(int) acc_device_*' cast to avoid the C++ 'acc_on_device' wrapper.  */
+    x[i] = acc_on_device((int) acc_device_not_host) ? fact_nohost(i) : 0;
+
+  for (int i = 0; i < N; ++i)
+    {
+      if (acc_get_device_type() == acc_device_host)
+	assert(x[i] == 0);
+      else
+	assert(x[i] == fact(i));
+    }
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-nohost-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-nohost-2.c
new file mode 100644
index 00000000000..4d081f269eb
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-nohost-2.c
@@ -0,0 +1,39 @@
+/* Test 'nohost' clause via 'weak'.
+
+   { dg-require-effective-target weak_undefined }
+
+   When the OpenACC 'routine' with 'nohost' clauses gets discarded, the weak symbol then resolves to 'NULL'.
+*/
+
+/* { dg-additional-sources routine-nohost-2_2.c } */
+
+/* { dg-additional-options "-fno-inline" } for stable results regarding OpenACC 'routine'.  */
+
+#include <assert.h>
+#include <openacc.h>
+
+#pragma acc routine //nohost
+__attribute__((weak))
+extern int f1(int);
+
+int main()
+{
+  int x = -10;
+
+#pragma acc serial copy(x)
+  /* { dg-warning {using vector_length \(32\), ignoring 1} "" { target openacc_nvidia_accel_selected } .-1 } */
+  {
+    if (f1)
+      x = f1(x);
+    else
+      x = 0;
+
+  }
+
+  if (acc_get_device_type() == acc_device_host)
+    assert(x == 0);
+  else
+    assert(x == -20);
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-nohost-2_2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-nohost-2_2.c
new file mode 100644
index 00000000000..60295459792
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-nohost-2_2.c
@@ -0,0 +1,18 @@
+/* { dg-skip-if "" { *-*-* } } */
+
+#pragma acc routine nohost
+int f1(int x)
+{
+  /* Make sure this fails host compilation.  */
+#if defined ACC_DEVICE_TYPE_host
+  asm ("IT'S A TRAP");
+#elif defined ACC_DEVICE_TYPE_nvidia
+  asm ("{\n\t  .reg .u32 %tid_x;\n\t  mov.u32 %tid_x, %tid.x;\n\t}");
+#elif defined ACC_DEVICE_TYPE_radeon
+  asm ("s_nop 0");
+#else
+# error Not ported to this ACC_DEVICE_TYPE
+#endif
+
+  return 2 * x;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/routine-nohost-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/routine-nohost-1.f90
new file mode 100644
index 00000000000..cd5bddc8685
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/routine-nohost-1.f90
@@ -0,0 +1,63 @@
+! Test 'nohost' clause via 'acc_on_device'.
+
+! { dg-do run }
+
+! With optimizations disabled, we currently don't expect that 'acc_on_device' "evaluates at compile time to a constant".
+! { dg-skip-if "TODO PR82391" { *-*-* } { "-O0" } }
+
+! { dg-additional-options "-fdump-tree-oaccdevlow" }
+
+program main
+  use openacc
+  implicit none
+  integer, parameter :: n = 10
+  integer :: a(n), i
+  integer, external :: fact_nohost
+  !$acc routine (fact_nohost)
+  integer, external :: fact
+
+  !$acc parallel loop
+  do i = 1, n
+     if (acc_on_device(acc_device_not_host)) then
+        a(i) = fact_nohost(i)
+     else
+        a(i) = 0
+     end if
+  end do
+  !$acc end parallel loop
+
+  do i = 1, n
+     if (acc_get_device_type() .eq. acc_device_host) then
+        if (a(i) .ne. 0) stop 10 + i
+     else
+        if (a(i) .ne. fact(i)) stop 20 + i
+     end if
+  end do
+end program main
+
+recursive function fact(x) result(res)
+  implicit none
+  !$acc routine (fact)
+  integer, intent(in) :: x
+  integer :: res
+
+  if (x < 1) then
+     res = 1
+  else
+     res = x * fact(x - 1)
+  end if
+end function fact
+
+function fact_nohost(x) result(res)
+  use openacc
+  implicit none
+  !$acc routine (fact_nohost) nohost
+  integer, intent(in) :: x
+  integer :: res
+  integer, external :: fact
+
+  res = fact(x)
+end function fact_nohost
+! { dg-final { scan-tree-dump-times {(?n)^OpenACC routine 'fact_nohost' has 'nohost' clause\.$} 1 oaccdevlow { target { ! offloading_enabled } } } }
+! { dg-final { scan-tree-dump-times {(?n)^OpenACC routine 'fact_nohost_' has 'nohost' clause\.$} 1 oaccdevlow { target offloading_enabled } } }
+!TODO See PR101551 for 'offloading_enabled' differences.
-- 
2.30.2