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>
This commit is contained in:
Thomas Schwinge
parent
6099b9cc8c
commit
a61f6afbee
|
|
@ -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,
|
||||
|
|
|
|||
|
|
@ -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
|
||||
|
|
|
|||
|
|
@ -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;
|
||||
|
||||
|
|
|
|||
|
|
@ -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
|
||||
|
|
|
|||
|
|
@ -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 ();
|
||||
|
|
|
|||
|
|
@ -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:
|
||||
|
|
|
|||
|
|
@ -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);
|
||||
|
|
|
|||
|
|
@ -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;
|
||||
|
|
|
|||
|
|
@ -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",
|
||||
|
|
|
|||
|
|
@ -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. */
|
||||
|
|
|
|||
|
|
@ -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);
|
||||
|
|
|
|||
|
|
@ -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);
|
||||
}
|
||||
|
|
|
|||
|
|
@ -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 ();
|
||||
}
|
||||
|
|
|
|||
|
|
@ -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;
|
||||
|
||||
|
|
|
|||
|
|
@ -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 ();
|
||||
}
|
||||
|
|
|
|||
|
|
@ -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. */
|
||||
|
|
|
|||
|
|
@ -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" } } */
|
||||
|
|
@ -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" } } */
|
||||
|
|
|
|||
|
|
@ -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);
|
||||
|
|
|
|||
|
|
@ -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 } } */
|
||||
|
|
@ -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} } */
|
||||
|
|
@ -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. */
|
||||
|
|
|
|||
|
|
@ -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.
|
||||
|
|
@ -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.
|
||||
|
|
|
|||
|
|
@ -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
|
||||
|
|
|
|||
|
|
@ -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
|
||||
|
|
|
|||
|
|
@ -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
|
||||
|
|
|
|||
|
|
@ -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
|
||||
|
|
|
|||
|
|
@ -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
|
||||
|
|
|
|||
|
|
@ -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
|
||||
|
|
|
|||
|
|
@ -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
|
||||
|
|
|
|||
|
|
@ -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
|
||||
|
|
|
|||
|
|
@ -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
|
||||
|
|
|
|||
|
|
@ -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
|
||||
|
|
|
|||
|
|
@ -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 ();
|
||||
|
|
|
|||
|
|
@ -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,
|
||||
|
|
|
|||
|
|
@ -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:
|
||||
|
|
|
|||
|
|
@ -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;
|
||||
}
|
||||
|
|
@ -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;
|
||||
}
|
||||
|
|
@ -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;
|
||||
}
|
||||
|
|
@ -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.
|
||||
Loading…
Reference in New Issue