From: Julian Brown Date: Thu, 7 Feb 2019 01:46:25 +0000 (-0800) Subject: Add support for OpenACC routine nohost clause X-Git-Url: http://git.ipfire.org/cgi-bin/gitweb.cgi?a=commitdiff_plain;h=c59ce21469a565983b5d6d5addcb28b4a9ae4746;p=thirdparty%2Fgcc.git Add support for OpenACC routine nohost clause 2018-10-02 Thomas Schwinge Cesar Philippidis 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-6.f90: Likewise. --- diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp index fe70dd7d1bfe..b3886838368d 100644 --- a/gcc/ChangeLog.omp +++ b/gcc/ChangeLog.omp @@ -1,3 +1,19 @@ +2018-10-02 Thomas Schwinge + Cesar Philippidis + + * 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. + 2020-04-19 Chung-Lin Tang PR other/76739 diff --git a/gcc/c-family/ChangeLog.omp b/gcc/c-family/ChangeLog.omp new file mode 100644 index 000000000000..0012a30a7b02 --- /dev/null +++ b/gcc/c-family/ChangeLog.omp @@ -0,0 +1,4 @@ +2018-10-02 Thomas Schwinge + Cesar Philippidis + + * c-pragma.h (pragma_omp_clause): Add PRAGMA_OACC_CLAUSE_NOHST. diff --git a/gcc/c-family/c-pragma.h b/gcc/c-family/c-pragma.h index 6c34ffa5be43..3aeca1d0e990 100644 --- a/gcc/c-family/c-pragma.h +++ b/gcc/c-family/c-pragma.h @@ -159,6 +159,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/ChangeLog.omp b/gcc/c/ChangeLog.omp index 523e82a4bdcf..9ed0068df651 100644 --- a/gcc/c/ChangeLog.omp +++ b/gcc/c/ChangeLog.omp @@ -1,3 +1,11 @@ +2018-10-02 Thomas Schwinge + Cesar Philippidis + + * c-parser.c (c_parser_omp_clause_name): Handle "nohost". + (c_parser_oacc_all_clauses): Handle PRAGMA_OACC_CLAUSE_NOHOST. + (c_parser_oacc_routine): Update. + * c-typeck.c (c_finish_omp_clauses): Handle OMP_CLAUSE_NOHOST. + 2020-04-19 Chung-Lin Tang PR other/76739 diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c index 5cdeb21a4580..7e138b85ec0b 100644 --- a/gcc/c/c-parser.c +++ b/gcc/c/c-parser.c @@ -12707,6 +12707,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)) @@ -16143,6 +16145,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, @@ -17070,7 +17077,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 96f5f777cae1..999f6ab1b4e2 100644 --- a/gcc/c/c-typeck.c +++ b/gcc/c/c-typeck.c @@ -15082,6 +15082,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/ChangeLog.omp b/gcc/cp/ChangeLog.omp index 6b63998284d1..881083d3d15c 100644 --- a/gcc/cp/ChangeLog.omp +++ b/gcc/cp/ChangeLog.omp @@ -1,7 +1,15 @@ +2018-10-02 Thomas Schwinge + Cesar Philippidis + + * parser.c (cp_parser_omp_clause_name): Handle "nohost". + (cp_parser_oacc_all_clauses): Handle PRAGMA_OACC_CLAUSE_NOHOST, + (cp_parser_oacc_routine): Update. + * pt.c (tsubst_omp_clauses): Handle OMP_CLAUSE_NOHOST. + * semantics.c (finish_omp_clauses): Handle OMP_CLAUSE_NOHOST. + 2020-04-19 Chung-Lin Tang PR other/76739 - * semantics.c (handle_omp_array_sections_1): Add 'bool &non_contiguous' parameter, adjust recursive call site, add cases for allowing pointer based multi-dimensional arrays for OpenACC. diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c index 99eccf0c5e43..51cdcc8c840c 100644 --- a/gcc/cp/parser.c +++ b/gcc/cp/parser.c @@ -35154,6 +35154,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)) @@ -38268,6 +38270,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"; @@ -44167,8 +44174,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 19fdafa4c43c..bd04e856181e 100644 --- a/gcc/cp/pt.c +++ b/gcc/cp/pt.c @@ -17597,6 +17597,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 aefe3f97eb07..9a3ecc1390e6 100644 --- a/gcc/cp/semantics.c +++ b/gcc/cp/semantics.c @@ -8156,6 +8156,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/ChangeLog.omp b/gcc/fortran/ChangeLog.omp index 38e2f39512d6..24f67c1568da 100644 --- a/gcc/fortran/ChangeLog.omp +++ b/gcc/fortran/ChangeLog.omp @@ -1,3 +1,16 @@ +2018-10-02 Thomas Schwinge + Cesar Philippidis + + * 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. + 2020-04-19 Chung-Lin Tang PR other/76739 diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h index 7935aca23db2..ce1252f98b3e 100644 --- a/gcc/fortran/gfortran.h +++ b/gcc/fortran/gfortran.h @@ -938,6 +938,7 @@ typedef struct /* OpenACC 'routine' directive's level of parallelism. */ ENUM_BITFIELD (oacc_routine_lop) oacc_routine_lop:3; + unsigned oacc_function_nohost:1; /* Attributes set by compiler extensions (!GCC$ ATTRIBUTES). */ unsigned ext_attr:EXT_ATTR_NUM; @@ -1435,7 +1436,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 par_auto:1, gang_static:1; - unsigned if_present:1, finalize:1; + unsigned if_present:1, finalize:1, nohost:1; locus loc; } diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c index 1f1920cf0e12..068cf79307eb 100644 --- a/gcc/fortran/openmp.c +++ b/gcc/fortran/openmp.c @@ -875,6 +875,7 @@ enum omp_mask2 OMP_CLAUSE_IF_PRESENT, OMP_CLAUSE_FINALIZE, OMP_CLAUSE_ATTACH, + OMP_CLAUSE_NOHOST, /* This must come last. */ OMP_MASK2_LAST }; @@ -1765,6 +1766,13 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, &c->lists[OMP_LIST_NONTEMPORAL], true) == MATCH_YES) 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 @@ -2283,7 +2291,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 @@ -2749,6 +2758,8 @@ 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_function_nohost + = c ? c->nohost : false; } else /* Something has gone wrong, possibly a syntax error. */ diff --git a/gcc/fortran/trans-decl.c b/gcc/fortran/trans-decl.c index cc9d85543ca6..952166334775 100644 --- a/gcc/fortran/trans-decl.c +++ b/gcc/fortran/trans-decl.c @@ -1496,8 +1496,14 @@ add_attributes_to_decl (symbol_attribute sym_attr, tree list) || sym_attr.oacc_declare_copyin || sym_attr.oacc_declare_deviceptr || sym_attr.oacc_declare_device_resident) - list = tree_cons (get_identifier ("omp declare target"), - clauses, list); + { + list = tree_cons (get_identifier ("omp declare target"), + clauses, list); + tree c = clauses; + 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); + } return list; } diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c index 349df1cc3468..c5c8e3f0cee6 100644 --- a/gcc/fortran/trans-openmp.c +++ b/gcc/fortran/trans-openmp.c @@ -1607,7 +1607,6 @@ gfc_omp_firstprivatize_type_sizes (struct gimplify_omp_ctx *ctx, tree type) } } - static inline tree gfc_trans_add_clause (tree node, tree tail) { @@ -3918,6 +3917,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 00f56d3fbdbb..fa7a002e5845 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -10044,6 +10044,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, } break; + case OMP_CLAUSE_NOHOST: default: gcc_unreachable (); } @@ -10969,6 +10970,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-low.c b/gcc/omp-low.c index 6d3602154b3e..6de26fd8e631 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -1787,6 +1787,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 (); @@ -1973,6 +1974,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) case OMP_CLAUSE__CONDTEMP_: 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 ba0937fba940..35cebe5db1a5 100644 --- a/gcc/omp-offload.c +++ b/gcc/omp-offload.c @@ -1745,6 +1745,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). */ @@ -1752,12 +1771,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) { @@ -1780,7 +1806,7 @@ execute_oacc_device_lower () bool is_oacc_parallel_kernels_gang_single = (lookup_attribute ("oacc parallel_kernels_gang_single", DECL_ATTRIBUTES (current_function_decl)) != NULL); - int fn_level = oacc_fn_attrib_level (attrs); + int fn_level = oacc_fn_attrib_level (attr); bool is_oacc_routine = (fn_level >= 0); gcc_checking_assert (is_oacc_parallel + is_oacc_kernels @@ -1825,11 +1851,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 (); + fn_level = oacc_fn_attrib_level (attr); unsigned outer_mask = 0; if (is_oacc_routine) @@ -1845,7 +1872,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/ChangeLog.omp b/gcc/testsuite/ChangeLog.omp index 64bb0cb2e5ce..0854ba050ba1 100644 --- a/gcc/testsuite/ChangeLog.omp +++ b/gcc/testsuite/ChangeLog.omp @@ -1,3 +1,12 @@ +2018-10-02 Thomas Schwinge + Cesar Philippidis + + * c-c++-common/goacc/routine-1.c: Adjust test. + * 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. + 2020-04-19 Chung-Lin Tang PR other/76739 diff --git a/gcc/testsuite/c-c++-common/goacc/routine-1.c b/gcc/testsuite/c-c++-common/goacc/routine-1.c index a75692246b64..db1322e11ca8 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 be1510a369ca..0e07435b571a 100644 --- a/gcc/testsuite/c-c++-common/goacc/routine-2.c +++ b/gcc/testsuite/c-c++-common/goacc/routine-2.c @@ -1,3 +1,23 @@ /* Test invalid use of the OpenACC 'routine' directive. */ +#pragma acc routine gang worker /* { dg-error "conflicting level" } */ +void gang (void) +{ +} + +#pragma acc routine worker vector /* { dg-error "conflicting level" } */ +void worker (void) +{ +} + +#pragma acc routine vector seq /* { dg-error "conflicting level" } */ +void vector (void) +{ +} + +#pragma acc routine seq gang /* { dg-error "conflicting level" } */ +void seq (void) +{ +} + #pragma acc routine (nothing) gang /* { dg-error "not been declared" } */ 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 000000000000..9baa56cb2060 --- /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 ea7c9bf73939..c82493321cb6 100644 --- a/gcc/testsuite/g++.dg/goacc/routine-2.C +++ b/gcc/testsuite/g++.dg/goacc/routine-2.C @@ -2,15 +2,8 @@ template 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 -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 000000000000..b295a4fcc59c --- /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 07ddf91a2303..8d1002c1e06f 100644 --- a/gcc/tree-core.h +++ b/gcc/tree-core.h @@ -495,6 +495,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 cea917a4d58a..1ce25f11dd45 100644 --- a/gcc/tree-nested.c +++ b/gcc/tree-nested.c @@ -1488,6 +1488,8 @@ convert_nonlocal_omp_clauses (tree *pclauses, struct walk_stmt_info *wi) case OMP_CLAUSE__SCANTEMP_: 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_: @@ -2268,6 +2270,8 @@ convert_local_omp_clauses (tree *pclauses, struct walk_stmt_info *wi) case OMP_CLAUSE__SCANTEMP_: 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 3a2f01bfe603..f64413ff3246 100644 --- a/gcc/tree-pretty-print.c +++ b/gcc/tree-pretty-print.c @@ -1310,6 +1310,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_IF_PRESENT: pp_string (pp, "if_present"); diff --git a/gcc/tree.c b/gcc/tree.c index e4e74ac8afc3..efcdafee057c 100644 --- a/gcc/tree.c +++ b/gcc/tree.c @@ -358,6 +358,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 */ 0, /* OMP_CLAUSE_IF_PRESENT */ 0, /* OMP_CLAUSE_FINALIZE */ @@ -444,6 +445,7 @@ const char * const omp_clause_code_name[] = "num_gangs", "num_workers", "vector_length", + "nohost", "tile", "if_present", "finalize", @@ -12292,6 +12294,7 @@ walk_tree_1 (tree *tp, walk_tree_fn func, void *data, case OMP_CLAUSE_BIND: 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/ChangeLog.omp b/libgomp/ChangeLog.omp index 5234168e4b1e..15faa26336bf 100644 --- a/libgomp/ChangeLog.omp +++ b/libgomp/ChangeLog.omp @@ -1,3 +1,12 @@ +2018-10-02 Thomas Schwinge + Cesar Philippidis + + * 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-6.f90: Likewise. + 2020-04-19 Chung-Lin Tang PR other/76739 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 000000000000..2cdd6bf459c9 --- /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 + +#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 000000000000..365af9319bdc --- /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 000000000000..1bae09c2a9ab --- /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 +