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>
diff --git a/gcc/c-family/c-pragma.h b/gcc/c-family/c-pragma.h
index e4fd3c9..c5d11ce 100644
--- a/gcc/c-family/c-pragma.h
+++ b/gcc/c-family/c-pragma.h
@@ -160,6 +160,7 @@
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 9a56e0c..92d22d1 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -12744,6 +12744,8 @@
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 @@
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 @@
( (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 4f7ed67..5d6565b 100644
--- a/gcc/c/c-typeck.c
+++ b/gcc/c/c-typeck.c
@@ -15168,6 +15168,7 @@
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 821ce17..45216f0 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -35656,6 +35656,8 @@
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 @@
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 @@
( (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 94ca3bc..b396ddd 100644
--- a/gcc/cp/pt.c
+++ b/gcc/cp/pt.c
@@ -17479,6 +17479,7 @@
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 331daf8..f64b084 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -8267,6 +8267,7 @@
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 26841ee..8e4a101 100644
--- a/gcc/fortran/dump-parse-tree.c
+++ b/gcc/fortran/dump-parse-tree.c
@@ -926,6 +926,8 @@
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 f4a50d7..921aed9 100644
--- a/gcc/fortran/gfortran.h
+++ b/gcc/fortran/gfortran.h
@@ -947,6 +947,7 @@
/* 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 @@
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 321d3256..1804066 100644
--- a/gcc/fortran/module.c
+++ b/gcc/fortran/module.c
@@ -2088,6 +2088,7 @@
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 @@
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 @@
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 @@
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 357a1e1..520a435 100644
--- a/gcc/fortran/openmp.c
+++ b/gcc/fortran/openmp.c
@@ -880,6 +880,7 @@
OMP_CLAUSE_IF_PRESENT,
OMP_CLAUSE_FINALIZE,
OMP_CLAUSE_ATTACH,
+ OMP_CLAUSE_NOHOST,
/* This must come last. */
OMP_MASK2_LAST
};
@@ -2083,6 +2084,13 @@
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 @@
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_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_error ("Multiple loop axes specified for routine at %C");
goto cleanup;
}
+ nohost = c ? c->nohost : false;
if (isym != NULL)
{
@@ -3024,6 +3035,13 @@
" 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 @@
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 @@
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 @@
/* 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 @@
&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 a73ce8a..bf8783a 100644
--- a/gcc/fortran/trans-decl.c
+++ b/gcc/fortran/trans-decl.c
@@ -1473,6 +1473,14 @@
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 ace4faf..ac3f5f3 100644
--- a/gcc/fortran/trans-openmp.c
+++ b/gcc/fortran/trans-openmp.c
@@ -4297,6 +4297,8 @@
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 5d43f76..21ff32e 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -10310,6 +10310,7 @@
}
break;
+ case OMP_CLAUSE_NOHOST:
default:
gcc_unreachable ();
}
@@ -11247,6 +11248,7 @@
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 a1bb9d8..b46a537 100644
--- a/gcc/omp-general.c
+++ b/gcc/omp-general.c
@@ -2576,6 +2576,7 @@
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 @@
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 @@
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 @@
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 @@
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 e7049c8..2f735bc 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -1683,6 +1683,7 @@
break;
case OMP_CLAUSE__CACHE_:
+ case OMP_CLAUSE_NOHOST:
default:
gcc_unreachable ();
}
@@ -1869,6 +1870,7 @@
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 0320ea6..bfbb011 100644
--- a/gcc/omp-offload.c
+++ b/gcc/omp-offload.c
@@ -1981,6 +1981,42 @@
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 0000000..a58482f
--- /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 81fe369..cc0ba2b 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 @@
/* 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 be1510a..3bf33e8 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 0000000..c892741
--- /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 0000000..d9acb80
--- /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 51a3f54..f34fcfe 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 @@
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 0000000..0e06fb9
--- /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 52cc870..92d3243 100644
--- a/gcc/testsuite/gfortran.dg/goacc/classify-routine.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/classify-routine.f95
@@ -29,5 +29,12 @@
! 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 97d92c3..31233b3 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 @@
!$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 @@
!$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 f1e2aa3..3cd543e 100644
--- a/gcc/testsuite/gfortran.dg/goacc/routine-6.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/routine-6.f90
@@ -116,3 +116,13 @@
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 22524cc..0372e78 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 4e81f11..46eec3d 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 @@
!$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 eae0807..e796c1d 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 @@
!$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 a4ff549..80fe07a 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 835619c..10e1096 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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
! { 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 622a9d9..44ef453 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 54365ae..f332ed5 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 9391609..bfab988 100644
--- a/gcc/tree-core.h
+++ b/gcc/tree-core.h
@@ -508,7 +508,10 @@
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 9edd922..0c3fb02 100644
--- a/gcc/tree-nested.c
+++ b/gcc/tree-nested.c
@@ -1510,6 +1510,9 @@
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 @@
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 fde07df..7201bd7 100644
--- a/gcc/tree-pretty-print.c
+++ b/gcc/tree-pretty-print.c
@@ -1303,6 +1303,9 @@
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 bead1ac..c621f87 100644
--- a/gcc/tree.c
+++ b/gcc/tree.c
@@ -361,6 +361,7 @@
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 @@
"tile",
"if_present",
"finalize",
+ "nohost",
};
@@ -11165,6 +11167,7 @@
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 0000000..dc92727
--- /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 0000000..4d081f2
--- /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 0000000..6029545
--- /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 0000000..cd5bddc
--- /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.
