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=fc61d6296451096488253692ad08978e230ab7a3;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. (cherry picked from openacc-gcc-9-branch commit a6ec4f19ff50a91a7b0a8c3663e49c071b506016) --- diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp index 3095c0c79f99..d1e1ab54ea9e 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. + 2019-02-05 Julian Brown * omp-low.c (scan_sharing_clauses): Disallow dynamic (multidimensional) diff --git a/gcc/c-family/ChangeLog.omp b/gcc/c-family/ChangeLog.omp index 20f26d1bd8a4..0f348b2219a8 100644 --- a/gcc/c-family/ChangeLog.omp +++ b/gcc/c-family/ChangeLog.omp @@ -1,3 +1,10 @@ +2018-10-02 Thomas Schwinge + Cesar Philippidis + + * 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. + 2018-12-14 Julian Brown * c-pragma.h (pragma_omp_clause): Add PRAGMA_OACC_CLAUSE_ATTACH and diff --git a/gcc/c-family/c-attribs.c b/gcc/c-family/c-attribs.c index 93b210eed9ca..c4bbdf9d7399 100644 --- a/gcc/c-family/c-attribs.c +++ b/gcc/c-family/c-attribs.c @@ -437,7 +437,7 @@ const struct attribute_spec c_common_attribute_table[] = handle_omp_declare_simd_attribute, NULL }, { "simd", 0, 1, true, false, false, false, handle_simd_attribute, NULL }, - { "omp declare target", 0, 0, true, false, false, false, + { "omp declare target", 0, -1, true, false, false, false, handle_omp_declare_target_attribute, NULL }, { "omp declare target link", 0, 0, true, false, false, false, handle_omp_declare_target_attribute, NULL }, diff --git a/gcc/c-family/c-pragma.h b/gcc/c-family/c-pragma.h index 4746d41fe4e8..62bdaee80ebf 100644 --- a/gcc/c-family/c-pragma.h +++ b/gcc/c-family/c-pragma.h @@ -149,6 +149,7 @@ enum pragma_omp_clause { PRAGMA_OACC_CLAUSE_GANG, PRAGMA_OACC_CLAUSE_HOST, PRAGMA_OACC_CLAUSE_INDEPENDENT, + PRAGMA_OACC_CLAUSE_NOHOST, PRAGMA_OACC_CLAUSE_NUM_GANGS, PRAGMA_OACC_CLAUSE_NUM_WORKERS, PRAGMA_OACC_CLAUSE_PRESENT, diff --git a/gcc/c/ChangeLog.omp b/gcc/c/ChangeLog.omp index 4909c4595dbc..023607f1e4d9 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, c_finish_oacc_routine): Update. + * c-typeck.c (c_finish_omp_clauses): Handle OMP_CLAUSE_NOHOST. + 2018-10-16 Chung-Lin Tang * c-typeck.c (handle_omp_array_sections_1): Add 'bool &non_contiguous' diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c index b745e607497f..60e4ed5b63b3 100644 --- a/gcc/c/c-parser.c +++ b/gcc/c/c-parser.c @@ -11749,6 +11749,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)) @@ -14920,6 +14922,11 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask, clauses = c_parser_oacc_data_clause (parser, c_kind, clauses); c_name = "link"; break; + case PRAGMA_OACC_CLAUSE_NOHOST: + clauses = c_parser_oacc_simple_clause (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, @@ -15767,7 +15774,8 @@ c_parser_oacc_kernels_parallel (location_t loc, c_parser *parser, ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_GANG) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WORKER) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_SEQ) ) + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_SEQ) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NOHOST) ) /* Parse an OpenACC routine directive. For named directives, we apply immediately to the named function. For unnamed ones we then parse @@ -15920,7 +15928,7 @@ c_finish_oacc_routine (struct oacc_routine_data *data, tree fndecl, /* Add an "omp declare target" attribute. */ DECL_ATTRIBUTES (fndecl) = tree_cons (get_identifier ("omp declare target"), - NULL_TREE, DECL_ATTRIBUTES (fndecl)); + data->clauses, DECL_ATTRIBUTES (fndecl)); /* Remember that we've used this "#pragma acc routine". */ data->fndecl_seen = true; diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c index 0146f06ffcbe..481b55b065b9 100644 --- a/gcc/c/c-typeck.c +++ b/gcc/c/c-typeck.c @@ -14729,6 +14729,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 a01ffb4c7302..f9ba662d4aa6 100644 --- a/gcc/cp/ChangeLog.omp +++ b/gcc/cp/ChangeLog.omp @@ -1,3 +1,12 @@ +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, cp_finalize_oacc_routine): Update. + * pt.c (tsubst_omp_clauses): Handle OMP_CLAUSE_NOHOST. + * semantics.c (finish_omp_clauses): Handle OMP_CLAUSE_NOHOST. + 2018-10-16 Chung-Lin Tang * semantics.c (handle_omp_array_sections_1): Add 'bool &non_contiguous' diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c index ed90aa2f5cd8..418be26be4b6 100644 --- a/gcc/cp/parser.c +++ b/gcc/cp/parser.c @@ -32285,6 +32285,8 @@ cp_parser_omp_clause_name (cp_parser *parser) case 'n': if (!strcmp ("nogroup", p)) result = PRAGMA_OMP_CLAUSE_NOGROUP; + else if (!strcmp ("nohost", p)) + result = PRAGMA_OACC_CLAUSE_NOHOST; else if (!strcmp ("nontemporal", p)) result = PRAGMA_OMP_CLAUSE_NONTEMPORAL; else if (!strcmp ("notinbranch", p)) @@ -35126,6 +35128,11 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask, clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses); c_name = "link"; break; + case PRAGMA_OACC_CLAUSE_NOHOST: + clauses = cp_parser_oacc_simple_clause (here, OMP_CLAUSE_NOHOST, + clauses); + c_name = "nohost"; + break; case PRAGMA_OACC_CLAUSE_NUM_GANGS: code = OMP_CLAUSE_NUM_GANGS; c_name = "num_gangs"; @@ -40129,8 +40136,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 @@ -40347,7 +40354,7 @@ cp_finalize_oacc_routine (cp_parser *parser, tree fndecl, bool is_defn) /* Add an "omp declare target" attribute. */ DECL_ATTRIBUTES (fndecl) = tree_cons (get_identifier ("omp declare target"), - NULL_TREE, DECL_ATTRIBUTES (fndecl)); + parser->oacc_routine->clauses, DECL_ATTRIBUTES (fndecl)); /* Don't unset parser->oacc_routine here: we may still need it to diagnose wrong usage. But, remember that we've used this "#pragma acc diff --git a/gcc/cp/pt.c b/gcc/cp/pt.c index e682b6d51be3..85f54ee15d96 100644 --- a/gcc/cp/pt.c +++ b/gcc/cp/pt.c @@ -16413,6 +16413,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 0282f86fde91..15519e773e38 100644 --- a/gcc/cp/semantics.c +++ b/gcc/cp/semantics.c @@ -7643,6 +7643,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) case OMP_CLAUSE_SEQ: case OMP_CLAUSE_IF_PRESENT: case OMP_CLAUSE_FINALIZE: + case OMP_CLAUSE_NOHOST: break; case OMP_CLAUSE_TILE: diff --git a/gcc/fortran/ChangeLog.omp b/gcc/fortran/ChangeLog.omp index 86ed18bdad13..933e32e68294 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. + 2018-12-14 Julian Brown * gfortran.h (gfc_omp_map_op): Add OMP_MAP_ATTACH, OMP_MAP_DETACH. diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h index 7e039e50f6c4..5cd16e37fabc 100644 --- a/gcc/fortran/gfortran.h +++ b/gcc/fortran/gfortran.h @@ -916,6 +916,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; @@ -1360,7 +1361,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 b08593a2317c..bfceb4010e64 100644 --- a/gcc/fortran/openmp.c +++ b/gcc/fortran/openmp.c @@ -815,6 +815,7 @@ enum omp_mask2 OMP_CLAUSE_FINALIZE, OMP_CLAUSE_ATTACH, OMP_CLAUSE_DETACH, + OMP_CLAUSE_NOHOST, /* This must come last. */ OMP_MASK2_LAST }; @@ -1464,6 +1465,13 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, c->nogroup = needs_space = true; continue; } + if ((mask & OMP_CLAUSE_NOHOST) + && !c->nohost + && gfc_match ("nohost") == MATCH_YES) + { + c->nohost = true; + continue; + } if ((mask & OMP_CLAUSE_NOTINBRANCH) && !c->notinbranch && !c->inbranch @@ -2003,7 +2011,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 @@ -2445,6 +2454,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 a0e1f6aeea56..bde3ca3907f8 100644 --- a/gcc/fortran/trans-decl.c +++ b/gcc/fortran/trans-decl.c @@ -1404,8 +1404,12 @@ add_attributes_to_decl (symbol_attribute sym_attr, tree list) list = tree_cons (get_identifier ("omp declare target link"), NULL_TREE, list); else if (sym_attr.omp_declare_target) - list = tree_cons (get_identifier ("omp declare target"), - NULL_TREE, list); + { + tree c = NULL_TREE; + if (sym_attr.oacc_function_nohost) + c = build_omp_clause (UNKNOWN_LOCATION, OMP_CLAUSE_NOHOST); + list = tree_cons (get_identifier ("omp declare target"), c, list); + } if (sym_attr.oacc_routine_lop != OACC_ROUTINE_LOP_NONE) { diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c index 0fe426a4edba..6c30f0620ceb 100644 --- a/gcc/fortran/trans-openmp.c +++ b/gcc/fortran/trans-openmp.c @@ -1315,7 +1315,6 @@ gfc_omp_firstprivatize_type_sizes (struct gimplify_omp_ctx *ctx, tree type) } } - static inline tree gfc_trans_add_clause (tree node, tree tail) { @@ -3088,6 +3087,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 dc4531bc2d30..ed9ceb633456 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -9461,6 +9461,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, ctx->default_kind = OMP_CLAUSE_DEFAULT_KIND (c); break; + case OMP_CLAUSE_NOHOST: default: gcc_unreachable (); } @@ -10239,6 +10240,7 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, case OMP_CLAUSE_FINALIZE: break; + case OMP_CLAUSE_NOHOST: default: gcc_unreachable (); } diff --git a/gcc/omp-low.c b/gcc/omp-low.c index a6bf2200f378..c105542d265d 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -1589,6 +1589,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 (); @@ -1765,6 +1766,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) case OMP_CLAUSE_FINALIZE: break; + case OMP_CLAUSE_NOHOST: case OMP_CLAUSE__CACHE_: default: gcc_unreachable (); diff --git a/gcc/omp-offload.c b/gcc/omp-offload.c index 97ae47b31358..51cf7f81caf6 100644 --- a/gcc/omp-offload.c +++ b/gcc/omp-offload.c @@ -1490,6 +1490,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). */ @@ -1497,12 +1516,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) { @@ -1523,12 +1549,12 @@ execute_oacc_device_lower () if (is_oacc_kernels && !is_oacc_kernels_parallelized) { oacc_set_fn_attrib (current_function_decl, NULL, NULL); - attrs = oacc_get_fn_attrib (current_function_decl); + attr = oacc_get_fn_attrib (current_function_decl); } /* Discover, partition and process the loops. */ oacc_loop *loops = oacc_loop_discovery (); - int fn_level = oacc_fn_attrib_level (attrs); + int fn_level = oacc_fn_attrib_level (attr); if (dump_file) { @@ -1555,7 +1581,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 7e5117da871f..2a0d2f461580 100644 --- a/gcc/testsuite/ChangeLog.omp +++ b/gcc/testsuite/ChangeLog.omp @@ -1,3 +1,13 @@ +2018-10-02 Thomas Schwinge + Cesar Philippidis + + * 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. + 2019-02-05 Julian Brown * c-c++-common/goacc/deep-copy-multidim.c: Add test. diff --git a/gcc/testsuite/c-c++-common/goacc/classify-routine.c b/gcc/testsuite/c-c++-common/goacc/classify-routine.c index a723d2cdf513..0b9ba6ea69fc 100644 --- a/gcc/testsuite/c-c++-common/goacc/classify-routine.c +++ b/gcc/testsuite/c-c++-common/goacc/classify-routine.c @@ -22,10 +22,10 @@ void ROUTINE () } /* Check the offloaded function's attributes. - { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(omp declare target, oacc function \\(0 1, 1 0, 1 0\\)\\)\\)" 1 "ompexp" } } */ + { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(omp declare target \\(worker\\), oacc function \\(0 1, 1 0, 1 0\\)\\)\\)" 1 "ompexp" } } */ /* Check the offloaded function's classification and compute dimensions (will always be 1 x 1 x 1 for non-offloading compilation). { dg-final { scan-tree-dump-times "(?n)Function is OpenACC routine level 1" 1 "oaccdevlow" } } { dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccdevlow" } } - { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(0 1, 1 1, 1 1\\), omp declare target, oacc function \\(0 1, 1 0, 1 0\\)\\)\\)" 1 "oaccdevlow" } } */ + { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(0 1, 1 1, 1 1\\), omp declare target \\(worker\\), oacc function \\(0 1, 1 0, 1 0\\)\\)\\)" 1 "oaccdevlow" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/routine-1.c b/gcc/testsuite/c-c++-common/goacc/routine-1.c index 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 fc5eb11bb54d..d1ea61e33105 100644 --- a/gcc/testsuite/c-c++-common/goacc/routine-2.c +++ b/gcc/testsuite/c-c++-common/goacc/routine-2.c @@ -1,19 +1,19 @@ -#pragma acc routine gang worker /* { dg-error "multiple loop axes" } */ +#pragma acc routine gang worker /* { dg-error "conflicting level" } */ void gang (void) { } -#pragma acc routine worker vector /* { dg-error "multiple loop axes" } */ +#pragma acc routine worker vector /* { dg-error "conflicting level" } */ void worker (void) { } -#pragma acc routine vector seq /* { dg-error "multiple loop axes" } */ +#pragma acc routine vector seq /* { dg-error "conflicting level" } */ void vector (void) { } -#pragma acc routine seq gang /* { dg-error "multiple loop axes" } */ +#pragma acc routine seq gang /* { dg-error "conflicting level" } */ void seq (void) { } diff --git a/gcc/testsuite/c-c++-common/goacc/routine-nohost-1.c b/gcc/testsuite/c-c++-common/goacc/routine-nohost-1.c new file mode 100644 index 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 552196b1b426..5bcecc160bc0 100644 --- a/gcc/tree-core.h +++ b/gcc/tree-core.h @@ -461,6 +461,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 3fe23cc2b22a..76071cc3c201 100644 --- a/gcc/tree-nested.c +++ b/gcc/tree-nested.c @@ -1350,6 +1350,8 @@ convert_nonlocal_omp_clauses (tree *pclauses, struct walk_stmt_info *wi) case OMP_CLAUSE_FINALIZE: break; + /* OpenACC nohost clause is not yet handled here. */ + case OMP_CLAUSE_NOHOST: /* The following clause belongs to the OpenACC cache directive, which is discarded during gimplification. */ case OMP_CLAUSE__CACHE_: @@ -2077,6 +2079,8 @@ convert_local_omp_clauses (tree *pclauses, struct walk_stmt_info *wi) case OMP_CLAUSE_FINALIZE: break; + /* OpenACC nohost clauses is not yet handled here. */ + case OMP_CLAUSE_NOHOST: /* The following clause belongs to the OpenACC cache directive, which is discarded during gimplification. */ case OMP_CLAUSE__CACHE_: diff --git a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c index 589972ed90d6..415b523a0f41 100644 --- a/gcc/tree-pretty-print.c +++ b/gcc/tree-pretty-print.c @@ -1200,6 +1200,9 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags) spc, flags, false); pp_right_paren (pp); break; + case OMP_CLAUSE_NOHOST: + pp_string (pp, "nohost"); + break; case OMP_CLAUSE__GRIDDIM_: pp_string (pp, "_griddim_("); diff --git a/gcc/tree.c b/gcc/tree.c index 73102c4e75bd..7c891dcbf914 100644 --- a/gcc/tree.c +++ b/gcc/tree.c @@ -347,6 +347,7 @@ unsigned const char omp_clause_num_ops[] = 1, /* OMP_CLAUSE_NUM_GANGS */ 1, /* OMP_CLAUSE_NUM_WORKERS */ 1, /* OMP_CLAUSE_VECTOR_LENGTH */ + 0, /* OMP_CLAUSE_NOHOST */ 3, /* OMP_CLAUSE_TILE */ 2, /* OMP_CLAUSE__GRIDDIM_ */ 0, /* OMP_CLAUSE_IF_PRESENT */ @@ -424,6 +425,7 @@ const char * const omp_clause_code_name[] = "num_gangs", "num_workers", "vector_length", + "nohost", "tile", "_griddim_", "if_present", @@ -12323,6 +12325,7 @@ walk_tree_1 (tree *tp, walk_tree_fn func, void *data, case OMP_CLAUSE_DEFAULTMAP: case OMP_CLAUSE_AUTO: case OMP_CLAUSE_SEQ: + case OMP_CLAUSE_NOHOST: case OMP_CLAUSE_TILE: case OMP_CLAUSE__SIMT_: case OMP_CLAUSE_IF_PRESENT: diff --git a/libgomp/ChangeLog.omp b/libgomp/ChangeLog.omp index 990199a30bf8..f5c8a47a9c55 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. + 2018-10-16 Chung-Lin Tang * testsuite/libgomp.oacc-c-c++-common/da-1.c: New test. 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 +