* [patch,openacc] Add support for OpenACC routine nohost clause
@ 2018-10-02 14:12 Cesar Philippidis
2021-07-21 22:20 ` OpenACC 'nohost' clause Thomas Schwinge
0 siblings, 1 reply; 2+ messages in thread
From: Cesar Philippidis @ 2018-10-02 14:12 UTC (permalink / raw)
To: gcc-patches, Schwinge, Thomas
[-- Attachment #1: Type: text/plain, Size: 682 bytes --]
Attached is a patch that introduces support for the acc routine nohost
clause. Basically, if an acc routine function is marked as nohost, then
the compiler does not generate code for the host. It's kind of strange
to test for. Basically, we had to use acc_on_device at -O2 so that the
host references to the dead function get optimized away.
I believe that the nohost clause was added for acc routines to allow
offloaded acc code to call vendor libraries, such as cuBLAS, which are
only available for specific accelerators. I haven't seen it used much in
practice though.
Is this OK for trunk? I bootstrapped and regtested it for x86_64 Linux
with nvptx offloading.
Thanks
Cesar
[-- Attachment #2: 0004-OpenACC-bind-nohost-changes.patch --]
[-- Type: text/x-patch, Size: 28779 bytes --]
[OpenACC] Add support for OpenACC routine nohost clause
(was OpenACC bind, nohost changes)
2018-XX-YY Thomas Schwinge <thomas@codesourcery.com>
Cesar Philippidis <cesar@codesourcery.com>
gcc/
* tree-core.h (omp_clause_code): Add OMP_CLAUSE_NOHOST.
* tree.c (omp_clause_num_ops, omp_clause_code_name, walk_tree_1):
Update for these.
* tree-pretty-print.c (dump_omp_clause): Handle OMP_CLAUSE_NOHOST.
* gimplify.c (gimplify_scan_omp_clauses)
(gimplify_adjust_omp_clauses): Handle OMP_CLAUSE_NOHOST.
* tree-nested.c (convert_nonlocal_omp_clauses)
(convert_local_omp_clauses): Likewise.
* omp-low.c (scan_sharing_clauses): Likewise.
* omp-offload.c (maybe_discard_oacc_function): New function.
(execute_oacc_device_lower) [!ACCEL_COMPILER]: Handle OpenACC
nohost clauses.
gcc/c-family/
* c-attribs.c (c_common_attribute_table): Set min_len to -1 for
"omp declare target".
* c-pragma.h (pragma_omp_clause): Add PRAGMA_OACC_CLAUSE_NOHST.
gcc/c/
* c-parser.c (c_parser_omp_clause_name): Handle "nohost".
(c_parser_oacc_all_clauses): Handle PRAGMA_OACC_CLAUSE_NOHOST.
(c_parser_oacc_routine, c_finish_oacc_routine): Update.
* 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,
(cp_parser_oacc_routine, cp_finalize_oacc_routine): Update.
* pt.c (tsubst_omp_clauses): Handle OMP_CLAUSE_NOHOST.
* semantics.c (finish_omp_clauses): Handle OMP_CLAUSE_NOHOST.
gcc/fortran/
* gfortran.h (gfc_omp_clauses): Add nohost members.
* openmp.c (omp_mask2): Add OMP_CLAUSE_NOHOST.
(gfc_match_omp_clauses): Handle OMP_CLAUSE_NOHOST.
(gfc_match_oacc_routine): Set oacc_function_nohost when appropriate.
* gfortran.h (symbol_attribute): Add oacc_function_nohost member.
* trans-openmp.c (gfc_add_omp_offload_attributes): Use it to decide
whether to generate an OMP_CLAUSE_NOHOST clause.
(gfc_trans_omp_clauses_1): Unreachable code to generate an
OMP_CLAUSE_NOHOST clause.
gcc/testsuite/
* c-c++-common/goacc/classify-routine.c: Adjust test.
* c-c++-common/goacc/routine-1.c: Likewise.
* c-c++-common/goacc/routine-2.c: Likewise.
* c-c++-common/goacc/routine-nohost-1.c: New test.
* g++.dg/goacc/routine-2.C: Adjust test.
* gfortran.dg/goacc/pr72741.f90: New test.
libgomp/
* testsuite/libgomp.oacc-c-c++-common/routine-3.c: New test.
* testsuite/libgomp.oacc-c-c++-common/routine-nohost-1.c: New test.
* testsuite/libgomp.oacc-c-c++-common/routine-bind-nohost-1.c:
Update test.
* testsuite/libgomp.oacc-fortran/routine-8.f90: Likewise.
(cherry picked from gomp-4_0-branch r223007, r226192, r226259, r228915,
r228916, and r231423)
(cherry picked from gomp-4_0-branch r231973 and r231979)
(cherry picked from gomp-4_0-branch r238847)
---
gcc/c-family/c-attribs.c | 2 +-
gcc/c-family/c-pragma.h | 1 +
gcc/c/c-parser.c | 12 +++++-
gcc/c/c-typeck.c | 1 +
gcc/cp/parser.c | 13 +++++--
gcc/cp/pt.c | 1 +
gcc/cp/semantics.c | 1 +
gcc/fortran/gfortran.h | 3 +-
gcc/fortran/openmp.c | 29 +++++++-------
gcc/fortran/trans-openmp.c | 15 +++++++-
gcc/gimplify.c | 2 +
gcc/lto/lto.c | 1 +
gcc/omp-low.c | 2 +
gcc/omp-offload.c | 38 ++++++++++++++++---
.../c-c++-common/goacc/classify-routine.c | 4 +-
gcc/testsuite/c-c++-common/goacc/routine-1.c | 8 ++++
gcc/testsuite/c-c++-common/goacc/routine-2.c | 8 ++--
.../c-c++-common/goacc/routine-nohost-1.c | 28 ++++++++++++++
gcc/testsuite/g++.dg/goacc/routine-2.C | 9 +----
gcc/testsuite/gfortran.dg/goacc/pr72741.f90 | 30 +++++++++++++++
gcc/tree-core.h | 3 ++
gcc/tree-nested.c | 4 ++
gcc/tree-pretty-print.c | 3 ++
gcc/tree.c | 3 ++
.../libgomp.oacc-c-c++-common/routine-3.c | 33 ++++++++++++++++
.../routine-nohost-1.c | 18 +++++++++
.../libgomp.oacc-fortran/routine-6.f90 | 28 ++++++++++++++
27 files changed, 257 insertions(+), 43 deletions(-)
create mode 100644 gcc/testsuite/c-c++-common/goacc/routine-nohost-1.c
create mode 100644 gcc/testsuite/gfortran.dg/goacc/pr72741.f90
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/routine-3.c
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/routine-nohost-1.c
create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/routine-6.f90
diff --git a/gcc/c-family/c-attribs.c b/gcc/c-family/c-attribs.c
index 5454e09adbc..a8dc91ae93d 100644
--- a/gcc/c-family/c-attribs.c
+++ b/gcc/c-family/c-attribs.c
@@ -435,7 +435,7 @@ const struct attribute_spec c_common_attribute_table[] =
handle_omp_declare_simd_attribute, NULL },
{ "simd", 0, 1, true, false, false, false,
handle_simd_attribute, NULL },
- { "omp declare target", 0, 0, true, false, false, false,
+ { "omp declare target", 0, -1, true, false, false, false,
handle_omp_declare_target_attribute, NULL },
{ "omp declare target link", 0, 0, true, false, false, false,
handle_omp_declare_target_attribute, NULL },
diff --git a/gcc/c-family/c-pragma.h b/gcc/c-family/c-pragma.h
index b322547b11a..7ef281ce8af 100644
--- a/gcc/c-family/c-pragma.h
+++ b/gcc/c-family/c-pragma.h
@@ -142,6 +142,7 @@ enum pragma_omp_clause {
PRAGMA_OACC_CLAUSE_GANG,
PRAGMA_OACC_CLAUSE_HOST,
PRAGMA_OACC_CLAUSE_INDEPENDENT,
+ 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 3517cb783d9..187a2dec999 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -11457,6 +11457,8 @@ c_parser_omp_clause_name (c_parser *parser)
result = PRAGMA_OMP_CLAUSE_NOTINBRANCH;
else if (!strcmp ("nowait", p))
result = PRAGMA_OMP_CLAUSE_NOWAIT;
+ else if (!strcmp ("nohost", p))
+ result = PRAGMA_OACC_CLAUSE_NOHOST;
else if (!strcmp ("num_gangs", p))
result = PRAGMA_OACC_CLAUSE_NUM_GANGS;
else if (!strcmp ("num_tasks", p))
@@ -14127,6 +14129,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 = "link";
break;
+ case PRAGMA_OACC_CLAUSE_NOHOST:
+ clauses = c_parser_oacc_simple_clause (parser, 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,
@@ -14949,7 +14956,8 @@ c_parser_oacc_kernels_parallel (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
@@ -15110,7 +15118,7 @@ c_finish_oacc_routine (struct oacc_routine_data *data, tree fndecl,
/* Add an "omp declare target" attribute. */
DECL_ATTRIBUTES (fndecl)
= tree_cons (get_identifier ("omp declare target"),
- NULL_TREE, DECL_ATTRIBUTES (fndecl));
+ data->clauses, DECL_ATTRIBUTES (fndecl));
/* Remember that we've used this "#pragma acc routine". */
data->fndecl_seen = true;
diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c
index a5a7da0084c..48eaf26a672 100644
--- a/gcc/c/c-typeck.c
+++ b/gcc/c/c-typeck.c
@@ -14043,6 +14043,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
case OMP_CLAUSE_GANG:
case OMP_CLAUSE_WORKER:
case OMP_CLAUSE_VECTOR:
+ case OMP_CLAUSE_NOHOST:
case OMP_CLAUSE_TILE:
case OMP_CLAUSE_IF_PRESENT:
case OMP_CLAUSE_FINALIZE:
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index fa7ee7798ae..d56105ca177 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -31460,6 +31460,8 @@ cp_parser_omp_clause_name (cp_parser *parser)
case 'n':
if (!strcmp ("nogroup", p))
result = PRAGMA_OMP_CLAUSE_NOGROUP;
+ else if (!strcmp ("nohost", p))
+ result = PRAGMA_OACC_CLAUSE_NOHOST;
else if (!strcmp ("notinbranch", p))
result = PRAGMA_OMP_CLAUSE_NOTINBRANCH;
else if (!strcmp ("nowait", p))
@@ -33856,6 +33858,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 = "link";
break;
+ case PRAGMA_OACC_CLAUSE_NOHOST:
+ clauses = cp_parser_oacc_simple_clause (parser, OMP_CLAUSE_NOHOST,
+ clauses, here);
+ c_name = "nohost";
+ break;
case PRAGMA_OACC_CLAUSE_NUM_GANGS:
code = OMP_CLAUSE_NUM_GANGS;
c_name = "num_gangs";
@@ -38055,8 +38062,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
@@ -38282,7 +38289,7 @@ cp_finalize_oacc_routine (cp_parser *parser, tree fndecl, bool is_defn)
/* Add an "omp declare target" attribute. */
DECL_ATTRIBUTES (fndecl)
= tree_cons (get_identifier ("omp declare target"),
- NULL_TREE, DECL_ATTRIBUTES (fndecl));
+ parser->oacc_routine->clauses, DECL_ATTRIBUTES (fndecl));
/* Don't unset parser->oacc_routine here: we may still need it to
diagnose wrong usage. But, remember that we've used this "#pragma acc
diff --git a/gcc/cp/pt.c b/gcc/cp/pt.c
index b8b6545434b..7932dd2714b 100644
--- a/gcc/cp/pt.c
+++ b/gcc/cp/pt.c
@@ -16148,6 +16148,7 @@ tsubst_omp_clauses (tree clauses, enum c_omp_region_type ort,
case OMP_CLAUSE_IF_PRESENT:
case OMP_CLAUSE_FINALIZE:
break;
+ case OMP_CLAUSE_NOHOST:
default:
gcc_unreachable ();
}
diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c
index bf3c63a09a1..c4442200b74 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -7111,6 +7111,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_TILE:
diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h
index 657b5bb5a65..781dc2a7d17 100644
--- a/gcc/fortran/gfortran.h
+++ b/gcc/fortran/gfortran.h
@@ -913,6 +913,7 @@ typedef struct
/* This is an OpenACC acclerator function at level N - 1 */
ENUM_BITFIELD (oacc_function) oacc_function:3;
+ unsigned oacc_function_nohost:1;
/* Attributes set by compiler extensions (!GCC$ ATTRIBUTES). */
unsigned ext_attr:EXT_ATTR_NUM;
@@ -1354,7 +1355,7 @@ typedef struct gfc_omp_clauses
gfc_expr_list *tile_list;
unsigned async:1, gang:1, worker:1, vector:1, seq:1, independent:1;
unsigned wait:1, par_auto:1, gang_static:1;
- unsigned if_present:1, finalize:1;
+ unsigned if_present:1, finalize:1, nohost;
locus loc;
}
diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
index a046863445d..60ecaf54523 100644
--- a/gcc/fortran/openmp.c
+++ b/gcc/fortran/openmp.c
@@ -811,6 +811,7 @@ enum omp_mask2
OMP_CLAUSE_TILE,
OMP_CLAUSE_IF_PRESENT,
OMP_CLAUSE_FINALIZE,
+ OMP_CLAUSE_NOHOST,
/* This must come last. */
OMP_MASK2_LAST
};
@@ -1439,6 +1440,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 = true;
+ continue;
+ }
if ((mask & OMP_CLAUSE_NOTINBRANCH)
&& !c->notinbranch
&& !c->inbranch
@@ -1971,7 +1979,8 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
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
@@ -2348,22 +2357,12 @@ gfc_match_oacc_routine (void)
!= MATCH_YES))
return MATCH_ERROR;
- /* Scan for invalid routine geometry. */
dims = gfc_oacc_routine_dims (c);
if (dims == OACC_FUNCTION_NONE)
{
- gfc_error ("Multiple loop axes specified in !$ACC ROUTINE at %L",
- &old_loc);
-
- /* Don't abort early, because it's important to let the user
- know of any potential duplicate routine directives. */
- seen_error = true;
- }
- else if (dims == OACC_FUNCTION_AUTO)
- {
- gfc_warning (0, "Expected one of %<gang%>, %<worker%>, %<vector%> or "
- "%<seq%> clauses in !$ACC ROUTINE at %L", &old_loc);
- dims = OACC_FUNCTION_SEQ;
+ gfc_error ("Multiple loop axes specified for routine %C");
+ gfc_current_locus = old_loc;
+ return MATCH_ERROR;
}
if (sym != NULL)
@@ -2406,6 +2405,8 @@ gfc_match_oacc_routine (void)
goto cleanup;
gfc_current_ns->proc_name->attr.oacc_function = dims;
+ gfc_current_ns->proc_name->attr.oacc_function_nohost
+ = c ? c->nohost : false;
if (seen_error)
goto cleanup;
diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c
index 753272d84c2..0ee78d210d3 100644
--- a/gcc/fortran/trans-openmp.c
+++ b/gcc/fortran/trans-openmp.c
@@ -1290,8 +1290,12 @@ gfc_add_omp_offload_attributes (symbol_attribute sym_attr, tree list)
list = tree_cons (get_identifier ("omp declare target link"),
NULL_TREE, list);
else if (sym_attr.omp_declare_target)
- list = tree_cons (get_identifier ("omp declare target"),
- NULL_TREE, list);
+ {
+ tree c = NULL_TREE;
+ if (sym_attr.oacc_function_nohost)
+ c = build_omp_clause (UNKNOWN_LOCATION, OMP_CLAUSE_NOHOST);
+ list = tree_cons (get_identifier ("omp declare target"), c, list);
+ }
if (sym_attr.oacc_function != OACC_FUNCTION_NONE)
{
@@ -3053,6 +3057,13 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
OMP_CLAUSE_GANG_STATIC_EXPR (c) = arg;
}
}
+ if (clauses->nohost)
+ {
+ c = build_omp_clause (where.lb->location, OMP_CLAUSE_NOHOST);
+ omp_clauses = gfc_trans_add_clause (c, omp_clauses);
+ //TODO
+ gcc_unreachable();
+ }
return nreverse (omp_clauses);
}
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 509fc2f3f5b..32c3fac378c 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -8552,6 +8552,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
ctx->default_kind = OMP_CLAUSE_DEFAULT_KIND (c);
break;
+ case OMP_CLAUSE_NOHOST:
default:
gcc_unreachable ();
}
@@ -9310,6 +9311,7 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
case OMP_CLAUSE_FINALIZE:
break;
+ case OMP_CLAUSE_NOHOST:
default:
gcc_unreachable ();
}
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -1334,6 +1334,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
install_var_local (decl, ctx);
break;
+ case OMP_CLAUSE_NOHOST:
case OMP_CLAUSE__CACHE_:
default:
gcc_unreachable ();
@@ -1500,6 +1501,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
case OMP_CLAUSE_FINALIZE:
break;
+ case OMP_CLAUSE_NOHOST:
case OMP_CLAUSE__CACHE_:
default:
gcc_unreachable ();
diff --git a/gcc/omp-offload.c b/gcc/omp-offload.c
index 0abf0283c9e..c2ad801dc2a 100644
--- a/gcc/omp-offload.c
+++ b/gcc/omp-offload.c
@@ -1451,6 +1451,25 @@ default_goacc_reduction (gcall *call)
gsi_replace_with_seq (&gsi, seq, true);
}
+/* Determine whether DECL should be discarded in this offload
+ compilation. */
+
+static bool
+maybe_discard_oacc_function (tree decl)
+{
+ tree attr = lookup_attribute ("omp declare target", DECL_ATTRIBUTES (decl));
+
+ if (!attr)
+ return false;
+
+ enum omp_clause_code kind = OMP_CLAUSE_NOHOST;
+
+ if (omp_find_clause (TREE_VALUE (attr), kind))
+ return true;
+
+ return false;
+}
+
/* Main entry point for oacc transformations which run on the device
compiler after LTO, so we know what the target device is at this
point (including the host fallback). */
@@ -1458,12 +1477,19 @@ default_goacc_reduction (gcall *call)
static unsigned int
execute_oacc_device_lower ()
{
- tree attrs = oacc_get_fn_attrib (current_function_decl);
-
- if (!attrs)
+ tree attr = oacc_get_fn_attrib (current_function_decl);
+ if (!attr)
/* Not an offloaded function. */
return 0;
+ if (maybe_discard_oacc_function (current_function_decl))
+ {
+ if (dump_file)
+ fprintf (dump_file, "Discarding function\n");
+ TREE_ASM_WRITTEN (current_function_decl) = 1;
+ return TODO_discard_function;
+ }
+
/* Parse the default dim argument exactly once. */
if ((const void *)flag_openacc_dims != &flag_openacc_dims)
{
@@ -1484,12 +1510,12 @@ execute_oacc_device_lower ()
if (is_oacc_kernels && !is_oacc_kernels_parallelized)
{
oacc_set_fn_attrib (current_function_decl, NULL, NULL);
- attrs = oacc_get_fn_attrib (current_function_decl);
+ attr = oacc_get_fn_attrib (current_function_decl);
}
/* Discover, partition and process the loops. */
oacc_loop *loops = oacc_loop_discovery ();
- int fn_level = oacc_fn_attrib_level (attrs);
+ int fn_level = oacc_fn_attrib_level (attr);
if (dump_file)
{
@@ -1516,7 +1542,7 @@ execute_oacc_device_lower ()
}
int dims[GOMP_DIM_MAX];
- oacc_validate_dims (current_function_decl, attrs, dims, fn_level, used_mask);
+ oacc_validate_dims (current_function_decl, attr, dims, fn_level, used_mask);
if (dump_file)
{
diff --git a/gcc/testsuite/c-c++-common/goacc/classify-routine.c b/gcc/testsuite/c-c++-common/goacc/classify-routine.c
index fd89fc1ec66..f54c3942bbf 100644
--- a/gcc/testsuite/c-c++-common/goacc/classify-routine.c
+++ b/gcc/testsuite/c-c++-common/goacc/classify-routine.c
@@ -21,10 +21,10 @@ void ROUTINE ()
}
/* Check the offloaded function's attributes.
- { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(omp declare target, oacc function \\(0 1, 1 0, 1 0\\)\\)\\)" 1 "ompexp" } } */
+ { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(omp declare target \\(worker\\), oacc function \\(0 1, 1 0, 1 0\\)\\)\\)" 1 "ompexp" } } */
/* 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)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, 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" } } */
diff --git a/gcc/testsuite/c-c++-common/goacc/routine-1.c b/gcc/testsuite/c-c++-common/goacc/routine-1.c
index a75692246b6..db1322e11ca 100644
--- a/gcc/testsuite/c-c++-common/goacc/routine-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/routine-1.c
@@ -1,3 +1,4 @@
+/* Test valid use of clauses with routine. */
#pragma acc routine gang
void gang (void)
@@ -19,6 +20,11 @@ void seq (void)
{
}
+#pragma acc routine nohost
+void nohost (void)
+{
+}
+
int main ()
{
#pragma acc kernels num_gangs (32) num_workers (32) vector_length (32)
@@ -27,6 +33,7 @@ int main ()
worker ();
vector ();
seq ();
+ nohost ();
}
#pragma acc parallel num_gangs (32) num_workers (32) vector_length (32)
@@ -35,6 +42,7 @@ int main ()
worker ();
vector ();
seq ();
+ nohost ();
}
return 0;
diff --git a/gcc/testsuite/c-c++-common/goacc/routine-2.c b/gcc/testsuite/c-c++-common/goacc/routine-2.c
index fc5eb11bb54..d1ea61e3310 100644
--- a/gcc/testsuite/c-c++-common/goacc/routine-2.c
+++ b/gcc/testsuite/c-c++-common/goacc/routine-2.c
@@ -1,19 +1,19 @@
-#pragma acc routine gang worker /* { dg-error "multiple loop axes" } */
+#pragma acc routine gang worker /* { dg-error "conflicting level" } */
void gang (void)
{
}
-#pragma acc routine worker vector /* { dg-error "multiple loop axes" } */
+#pragma acc routine worker vector /* { dg-error "conflicting level" } */
void worker (void)
{
}
-#pragma acc routine vector seq /* { dg-error "multiple loop axes" } */
+#pragma acc routine vector seq /* { dg-error "conflicting level" } */
void vector (void)
{
}
-#pragma acc routine seq gang /* { dg-error "multiple loop axes" } */
+#pragma acc routine seq gang /* { dg-error "conflicting level" } */
void seq (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..9baa56cb206
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/routine-nohost-1.c
@@ -0,0 +1,28 @@
+/* Test the nohost clause for OpenACC routine directive. Exercising different
+ variants for declaring routines. */
+
+/* { dg-additional-options "-fdump-tree-oaccdevlow" } */
+
+#pragma acc routine nohost
+int THREE(void)
+{
+ return 3;
+}
+
+#pragma acc routine nohost
+extern void NOTHING(void);
+
+void NOTHING(void)
+{
+}
+
+extern float ADD(float, float);
+
+#pragma acc routine (ADD) nohost
+
+float ADD(float x, float y)
+{
+ return x + y;
+}
+
+/* { dg-final { scan-tree-dump-times "Discarding function" 3 "oaccdevlow" } } */
diff --git a/gcc/testsuite/g++.dg/goacc/routine-2.C b/gcc/testsuite/g++.dg/goacc/routine-2.C
index ea7c9bf7393..c82493321cb 100644
--- a/gcc/testsuite/g++.dg/goacc/routine-2.C
+++ b/gcc/testsuite/g++.dg/goacc/routine-2.C
@@ -2,15 +2,8 @@
template <typename T>
extern T one_d();
-#pragma acc routine (one_d) /* { dg-error "names a set of overloads" } */
+#pragma acc routine (one_d) nohost /* { dg-error "names a set of overloads" } */
-template <typename T>
-T
-one()
-{
- return 1;
-}
-#pragma acc routine (one) /* { dg-error "names a set of overloads" } */
int incr (int);
float incr (float);
diff --git a/gcc/testsuite/gfortran.dg/goacc/pr72741.f90 b/gcc/testsuite/gfortran.dg/goacc/pr72741.f90
new file mode 100644
index 00000000000..b295a4fcc59
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/pr72741.f90
@@ -0,0 +1,30 @@
+SUBROUTINE v_1
+ !$ACC ROUTINE VECTOR WORKER ! { dg-error "Multiple loop axes" }
+END SUBROUTINE v_1
+
+SUBROUTINE sub_1
+ IMPLICIT NONE
+ EXTERNAL :: g_1
+ !$ACC ROUTINE (g_1) GANG WORKER ! { dg-error "Multiple loop axes" }
+ !$ACC ROUTINE (ABORT) SEQ VECTOR ! { dg-error "Multiple loop axes" "" { xfail *-*-* } }
+! { dg-bogus "invalid function name abort" "" { xfail *-*-* } .-1 }
+
+ CALL v_1
+ CALL g_1
+ CALL ABORT
+END SUBROUTINE sub_1
+
+MODULE m_w_1
+ IMPLICIT NONE
+ EXTERNAL :: w_1
+ !$ACC ROUTINE (w_1) WORKER SEQ ! { dg-error "Multiple loop axes" }
+ !$ACC ROUTINE (ABORT) VECTOR GANG ! { dg-error "Multiple loop axes" "" { xfail *-*-* } }
+! { dg-bogus "invalid function name abort" "" { xfail *-*-* } .-1 }
+
+CONTAINS
+ SUBROUTINE sub_2
+ CALL v_1
+ CALL w_1
+ CALL ABORT
+ END SUBROUTINE sub_2
+END MODULE m_w_1
diff --git a/gcc/tree-core.h b/gcc/tree-core.h
index dee27f89dec..22fc9334a73 100644
--- a/gcc/tree-core.h
+++ b/gcc/tree-core.h
@@ -449,6 +449,9 @@ enum omp_clause_code {
/* OpenACC clause: vector_length (integer-expression). */
OMP_CLAUSE_VECTOR_LENGTH,
+ /* OpenACC clause: nohost. */
+ OMP_CLAUSE_NOHOST,
+
/* OpenACC clause: tile ( size-expr-list ). */
OMP_CLAUSE_TILE,
diff --git a/gcc/tree-nested.c b/gcc/tree-nested.c
index 4c8eda94f14..eddb91998c7 100644
--- a/gcc/tree-nested.c
+++ b/gcc/tree-nested.c
@@ -1345,6 +1345,8 @@ convert_nonlocal_omp_clauses (tree *pclauses, struct walk_stmt_info *wi)
case OMP_CLAUSE_FINALIZE:
break;
+ /* OpenACC nohost clause is not yet handled here. */
+ case OMP_CLAUSE_NOHOST:
/* The following clause belongs to the OpenACC cache directive, which
is discarded during gimplification. */
case OMP_CLAUSE__CACHE_:
@@ -2036,6 +2038,8 @@ convert_local_omp_clauses (tree *pclauses, struct walk_stmt_info *wi)
case OMP_CLAUSE_FINALIZE:
break;
+ /* OpenACC nohost clauses is not yet handled here. */
+ case OMP_CLAUSE_NOHOST:
/* The following clause belongs to the OpenACC cache directive, which
is discarded during gimplification. */
case OMP_CLAUSE__CACHE_:
diff --git a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c
index 2c089b11751..e6f32f12fd9 100644
--- a/gcc/tree-pretty-print.c
+++ b/gcc/tree-pretty-print.c
@@ -1043,6 +1043,9 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
spc, flags, false);
pp_right_paren (pp);
break;
+ case OMP_CLAUSE_NOHOST:
+ pp_string (pp, "nohost");
+ break;
case OMP_CLAUSE__GRIDDIM_:
pp_string (pp, "_griddim_(");
diff --git a/gcc/tree.c b/gcc/tree.c
index c3ac8f36d55..f953e08a0a8 100644
--- a/gcc/tree.c
+++ b/gcc/tree.c
@@ -341,6 +341,7 @@ unsigned const char omp_clause_num_ops[] =
1, /* OMP_CLAUSE_NUM_GANGS */
1, /* OMP_CLAUSE_NUM_WORKERS */
1, /* OMP_CLAUSE_VECTOR_LENGTH */
+ 0, /* OMP_CLAUSE_NOHOST */
3, /* OMP_CLAUSE_TILE */
2, /* OMP_CLAUSE__GRIDDIM_ */
0, /* OMP_CLAUSE_IF_PRESENT */
@@ -414,6 +415,7 @@ const char * const omp_clause_code_name[] =
"num_gangs",
"num_workers",
"vector_length",
+ "nohost",
"tile",
"_griddim_",
"if_present",
@@ -11672,6 +11674,7 @@ walk_tree_1 (tree *tp, walk_tree_fn func, void *data,
case OMP_CLAUSE_DEFAULTMAP:
case OMP_CLAUSE_AUTO:
case OMP_CLAUSE_SEQ:
+ case OMP_CLAUSE_NOHOST:
case OMP_CLAUSE_TILE:
case OMP_CLAUSE__SIMT_:
case OMP_CLAUSE_IF_PRESENT:
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-3.c
new file mode 100644
index 00000000000..2cdd6bf459c
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-3.c
@@ -0,0 +1,33 @@
+/* At -O0, we do get the expected "undefined reference to `foo'" link-time
+ error message (but the check needs to be done differently; compare to
+ routine-nohost-1.c), but for -O2 we don't; presumably because the function
+ gets inlined.
+ { dg-xfail-if "TODO" { *-*-* } { "-O0" } { "" } } */
+
+#include <stdlib.h>
+
+#pragma acc routine nohost
+int
+foo (int n)
+{
+ if (n == 0 || n == 1)
+ return 1;
+
+ return n * n;
+}
+
+int
+main()
+{
+ int a, n = 10;
+
+#pragma acc parallel copy (a, n)
+ {
+ a = foo (n);
+ }
+
+ if (a != n * n)
+ abort ();
+
+ return 0;
+}
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..365af9319bd
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-nohost-1.c
@@ -0,0 +1,18 @@
+/* { dg-do link } */
+
+extern int three (void);
+
+#pragma acc routine (three) nohost
+__attribute__((noinline))
+int three(void)
+{
+ return 3;
+}
+
+int main(void)
+{
+ return (three() == 3) ? 0 : 1;
+}
+
+/* Expecting link to fail; "undefined reference to `three'" (or similar).
+ { dg-excess-errors "" } */
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/routine-6.f90 b/libgomp/testsuite/libgomp.oacc-fortran/routine-6.f90
new file mode 100644
index 00000000000..1bae09c2a9a
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/routine-6.f90
@@ -0,0 +1,28 @@
+! { dg-do run }
+! { dg-xfail-if "TODO" { *-*-* } }
+
+program main
+ integer :: a, n
+
+ n = 10
+
+ !$acc parallel copy (a, n)
+ a = foo (n)
+ !$acc end parallel
+
+ if (a .ne. n * n) call abort
+
+contains
+
+function foo (n) result (rc)
+ !$acc routine nohost
+
+ integer, intent (in) :: n
+ integer :: rc
+
+ rc = n * n
+
+end function
+
+end program main
+
--
2.17.1
^ permalink raw reply [flat|nested] 2+ messages in thread
* OpenACC 'nohost' clause
2018-10-02 14:12 [patch,openacc] Add support for OpenACC routine nohost clause Cesar Philippidis
@ 2021-07-21 22:20 ` Thomas Schwinge
0 siblings, 0 replies; 2+ messages in thread
From: Thomas Schwinge @ 2021-07-21 22:20 UTC (permalink / raw)
To: gcc-patches, fortran
[-- Attachment #1: Type: text/plain, Size: 1587 bytes --]
Hi!
On 2018-10-02T07:11:43-0700, Cesar Philippidis <cesar@codesourcery.com> wrote:
> Attached is a patch that introduces support for the acc routine nohost
> clause. Basically, if an acc routine function is marked as nohost, then
> the compiler does not generate code for the host.
This is in particular useful in combination with the OpenACC 'bind'
clause and 'device_type' clause, which we don't have yet, so:
> It's kind of strange
> to test for. Basically, we had to use acc_on_device at -O2 so that the
> host references to the dead function get optimized away.
Additionally I figured out something using weak symbols.
> I believe that the nohost clause was added for acc routines to allow
> offloaded acc code to call vendor libraries, such as cuBLAS, which are
> only available for specific accelerators. I haven't seen it used much in
> practice though.
ACK.
> Is this OK for trunk?
After fixing the crucial issue to discard 'nohost' functions only for the
host but not also for all offload targets ;-) and considerably
improving/fixing the Fortran front end changes and boosting C/C++/Fortran
test coverage generally, I've now pushed "OpenACC 'nohost' clause" to
master branch in commit a61f6afbee370785cf091fe46e2e022748528307, see
attached.
Grüße
Thomas
-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955
[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #2: 0001-OpenACC-nohost-clause.patch --]
[-- Type: text/x-diff, Size: 84148 bytes --]
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
^ permalink raw reply [flat|nested] 2+ messages in thread
end of thread, other threads:[~2021-07-21 22:20 UTC | newest]
Thread overview: 2+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2018-10-02 14:12 [patch,openacc] Add support for OpenACC routine nohost clause Cesar Philippidis
2021-07-21 22:20 ` OpenACC 'nohost' clause Thomas Schwinge
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).