diff options
Diffstat (limited to 'gcc')
| -rw-r--r-- | gcc/ChangeLog | 23 | ||||
| -rw-r--r-- | gcc/c-family/ChangeLog | 8 | ||||
| -rw-r--r-- | gcc/c-family/c-pragma.c | 1 | ||||
| -rw-r--r-- | gcc/c-family/c-pragma.h | 5 | ||||
| -rw-r--r-- | gcc/c/ChangeLog | 12 | ||||
| -rw-r--r-- | gcc/c/c-parser.c | 176 | ||||
| -rw-r--r-- | gcc/cp/ChangeLog | 14 | ||||
| -rw-r--r-- | gcc/cp/parser.c | 172 | ||||
| -rw-r--r-- | gcc/cp/pt.c | 8 | ||||
| -rw-r--r-- | gcc/gimple-pretty-print.c | 3 | ||||
| -rw-r--r-- | gcc/gimple.h | 2 | ||||
| -rw-r--r-- | gcc/gimplify.c | 184 | ||||
| -rw-r--r-- | gcc/omp-builtins.def | 2 | ||||
| -rw-r--r-- | gcc/omp-low.c | 10 | ||||
| -rw-r--r-- | gcc/testsuite/ChangeLog | 6 | ||||
| -rw-r--r-- | gcc/testsuite/c-c++-common/goacc/declare-1.c | 83 | ||||
| -rw-r--r-- | gcc/testsuite/c-c++-common/goacc/declare-2.c | 79 | ||||
| -rw-r--r-- | gcc/tree-pretty-print.c | 6 |
18 files changed, 791 insertions, 3 deletions
diff --git a/gcc/ChangeLog b/gcc/ChangeLog index cb41838..5093ce0 100644 --- a/gcc/ChangeLog +++ b/gcc/ChangeLog @@ -1,3 +1,26 @@ +2015-11-12 James Norris <jnorris@codesourcery.com> + Joseph Myers <joseph@codesourcery.com> + + * gimple-pretty-print.c (dump_gimple_omp_target): Handle + GF_OMP_TARGET_KIND_OACC_DECLARE. + * gimple.h (enum gf_mask): Add GF_OMP_TARGET_KIND_OACC_DECLARE. + (is_gomple_omp_oacc): Handle GF_OMP_TARGET_KIND_OACC_DECLARE. + * gimplify.c (oacc_declare_returns): New. + (gimplify_bind_expr): Prepend 'exit' stmt to cleanup. + (device_resident_p): New function. + (oacc_default_clause): Handle device_resident clause. + (gimplify_oacc_declare_1, gimplify_oacc_declare): New functions. + (gimplify_expr): Handle OACC_DECLARE. + * omp-builtins.def (BUILT_IN_GOACC_DECLARE): New builtin. + * omp-low.c (expand_omp_target): Handle + GF_OMP_TARGET_KIND_OACC_DECLARE and BUILTIN_GOACC_DECLARE. + (build_omp_regions_1): Handlde GF_OMP_TARGET_KIND_OACC_DECLARE. + (lower_omp_target): Handle GF_OMP_TARGET_KIND_OACC_DECLARE, + GOMP_MAP_DEVICE_RESIDENT and GOMP_MAP_LINK. + (make_gimple_omp_edges): Handle GF_OMP_TARGET_KIND_OACC_DECLARE. + * tree-pretty-print.c (dump_omp_clause): Handle GOMP_MAP_LINK and + GOMP_MAP_DEVICE_RESIDENT. + 2015-11-12 Christophe Lyon <christophe.lyon@linaro.org> [ARM] Remove neon-testgen.ml and generated tests. diff --git a/gcc/c-family/ChangeLog b/gcc/c-family/ChangeLog index eb4d5bf..5611403 100644 --- a/gcc/c-family/ChangeLog +++ b/gcc/c-family/ChangeLog @@ -1,3 +1,11 @@ +2015-11-12 James Norris <jnorris@codesourcery.com> + Joseph Myers <joseph@codesourcery.com> + + * c-pragma.c (oacc_pragmas): Add entry for declare directive. + * c-pragma.h (enum pragma_kind): Add PRAGMA_OACC_DECLARE. + (enum pragma_omp_clause): Add PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT and + PRAGMA_OACC_CLAUSE_LINK. + 2015-11-11 Marek Polacek <polacek@redhat.com> PR c/68107 diff --git a/gcc/c-family/c-pragma.c b/gcc/c-family/c-pragma.c index f86ed38..12c3e75 100644 --- a/gcc/c-family/c-pragma.c +++ b/gcc/c-family/c-pragma.c @@ -1248,6 +1248,7 @@ static const struct omp_pragma_def oacc_pragmas[] = { { "atomic", PRAGMA_OACC_ATOMIC }, { "cache", PRAGMA_OACC_CACHE }, { "data", PRAGMA_OACC_DATA }, + { "declare", PRAGMA_OACC_DECLARE }, { "enter", PRAGMA_OACC_ENTER_DATA }, { "exit", PRAGMA_OACC_EXIT_DATA }, { "kernels", PRAGMA_OACC_KERNELS }, diff --git a/gcc/c-family/c-pragma.h b/gcc/c-family/c-pragma.h index afeceff..999ac67 100644 --- a/gcc/c-family/c-pragma.h +++ b/gcc/c-family/c-pragma.h @@ -30,6 +30,7 @@ enum pragma_kind { PRAGMA_OACC_ATOMIC, PRAGMA_OACC_CACHE, PRAGMA_OACC_DATA, + PRAGMA_OACC_DECLARE, PRAGMA_OACC_ENTER_DATA, PRAGMA_OACC_EXIT_DATA, PRAGMA_OACC_KERNELS, @@ -152,6 +153,7 @@ enum pragma_omp_clause { PRAGMA_OACC_CLAUSE_CREATE, PRAGMA_OACC_CLAUSE_DELETE, PRAGMA_OACC_CLAUSE_DEVICEPTR, + PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT, PRAGMA_OACC_CLAUSE_GANG, PRAGMA_OACC_CLAUSE_HOST, PRAGMA_OACC_CLAUSE_INDEPENDENT, @@ -176,7 +178,8 @@ enum pragma_omp_clause { PRAGMA_OACC_CLAUSE_FIRSTPRIVATE = PRAGMA_OMP_CLAUSE_FIRSTPRIVATE, PRAGMA_OACC_CLAUSE_IF = PRAGMA_OMP_CLAUSE_IF, PRAGMA_OACC_CLAUSE_PRIVATE = PRAGMA_OMP_CLAUSE_PRIVATE, - PRAGMA_OACC_CLAUSE_REDUCTION = PRAGMA_OMP_CLAUSE_REDUCTION + PRAGMA_OACC_CLAUSE_REDUCTION = PRAGMA_OMP_CLAUSE_REDUCTION, + PRAGMA_OACC_CLAUSE_LINK = PRAGMA_OMP_CLAUSE_LINK }; extern struct cpp_reader* parse_in; diff --git a/gcc/c/ChangeLog b/gcc/c/ChangeLog index 0191b45..02ae07e 100644 --- a/gcc/c/ChangeLog +++ b/gcc/c/ChangeLog @@ -1,3 +1,15 @@ +2015-11-12 James Norris <jnorris@codesourcery.com> + Joseph Myers <joseph@codesourcery.com> + + * c-parser.c (c_parser_pragma): Handle PRAGMA_OACC_DECLARE. + (c_parser_omp_clause_name): Handle 'device_resident' clause. + (c_parser_oacc_data_clause): Handle PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT + and PRAGMA_OMP_CLAUSE_LINK. + (c_parser_oacc_all_clauses): Handle PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT + and PRAGMA_OACC_CLAUSE_LINK. + (OACC_DECLARE_CLAUSE_MASK): New definition. + (c_parser_oacc_declare): New function. + 2015-11-12 Marek Polacek <polacek@redhat.com> PR c/67784 diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c index 8949825..c01d651 100644 --- a/gcc/c/c-parser.c +++ b/gcc/c/c-parser.c @@ -1228,6 +1228,7 @@ static vec<tree, va_gc> *c_parser_expr_list (c_parser *, bool, bool, vec<tree, va_gc> **, location_t *, tree *, vec<location_t> *, unsigned int * = NULL); +static void c_parser_oacc_declare (c_parser *); static void c_parser_oacc_enter_exit_data (c_parser *, bool); static void c_parser_oacc_update (c_parser *); static void c_parser_omp_construct (c_parser *); @@ -9729,6 +9730,10 @@ c_parser_pragma (c_parser *parser, enum pragma_context context) switch (id) { + case PRAGMA_OACC_DECLARE: + c_parser_oacc_declare (parser); + return false; + case PRAGMA_OACC_ENTER_DATA: c_parser_oacc_enter_exit_data (parser, true); return false; @@ -10018,6 +10023,8 @@ c_parser_omp_clause_name (c_parser *parser) result = PRAGMA_OMP_CLAUSE_DEVICE; else if (!strcmp ("deviceptr", p)) result = PRAGMA_OACC_CLAUSE_DEVICEPTR; + else if (!strcmp ("device_resident", p)) + result = PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT; else if (!strcmp ("dist_schedule", p)) result = PRAGMA_OMP_CLAUSE_DIST_SCHEDULE; break; @@ -10454,10 +10461,16 @@ c_parser_oacc_data_clause (c_parser *parser, pragma_omp_clause c_kind, case PRAGMA_OACC_CLAUSE_DEVICE: kind = GOMP_MAP_FORCE_TO; break; + case PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT: + kind = GOMP_MAP_DEVICE_RESIDENT; + break; case PRAGMA_OACC_CLAUSE_HOST: case PRAGMA_OACC_CLAUSE_SELF: kind = GOMP_MAP_FORCE_FROM; break; + case PRAGMA_OACC_CLAUSE_LINK: + kind = GOMP_MAP_LINK; + break; case PRAGMA_OACC_CLAUSE_PRESENT: kind = GOMP_MAP_FORCE_PRESENT; break; @@ -12739,6 +12752,10 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask, clauses = c_parser_oacc_data_clause_deviceptr (parser, clauses); c_name = "deviceptr"; break; + case PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT: + clauses = c_parser_oacc_data_clause (parser, c_kind, clauses); + c_name = "device_resident"; + break; case PRAGMA_OACC_CLAUSE_FIRSTPRIVATE: clauses = c_parser_omp_clause_firstprivate (parser, clauses); c_name = "firstprivate"; @@ -12761,6 +12778,10 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask, clauses); c_name = "independent"; break; + case PRAGMA_OACC_CLAUSE_LINK: + clauses = c_parser_oacc_data_clause (parser, c_kind, clauses); + c_name = "link"; + break; case PRAGMA_OACC_CLAUSE_NUM_GANGS: clauses = c_parser_omp_clause_num_gangs (parser, clauses); c_name = "num_gangs"; @@ -13218,6 +13239,161 @@ c_parser_oacc_data (location_t loc, c_parser *parser) } /* OpenACC 2.0: + # pragma acc declare oacc-data-clause[optseq] new-line +*/ + +#define OACC_DECLARE_CLAUSE_MASK \ + ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPY) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_LINK) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE) ) + +static void +c_parser_oacc_declare (c_parser *parser) +{ + location_t pragma_loc = c_parser_peek_token (parser)->location; + tree clauses, stmt, t, decl; + + bool error = false; + + c_parser_consume_pragma (parser); + + clauses = c_parser_oacc_all_clauses (parser, OACC_DECLARE_CLAUSE_MASK, + "#pragma acc declare"); + if (!clauses) + { + error_at (pragma_loc, + "no valid clauses specified in %<#pragma acc declare%>"); + return; + } + + for (t = clauses; t; t = OMP_CLAUSE_CHAIN (t)) + { + location_t loc = OMP_CLAUSE_LOCATION (t); + decl = OMP_CLAUSE_DECL (t); + if (!DECL_P (decl)) + { + error_at (loc, "array section in %<#pragma acc declare%>"); + error = true; + continue; + } + + switch (OMP_CLAUSE_MAP_KIND (t)) + { + case GOMP_MAP_FORCE_ALLOC: + case GOMP_MAP_FORCE_TO: + case GOMP_MAP_FORCE_DEVICEPTR: + case GOMP_MAP_DEVICE_RESIDENT: + break; + + case GOMP_MAP_POINTER: + /* Generated by c_finish_omp_clauses from array sections; + avoid spurious diagnostics. */ + break; + + case GOMP_MAP_LINK: + if (!global_bindings_p () + && (TREE_STATIC (decl) + || !DECL_EXTERNAL (decl))) + { + error_at (loc, + "%qD must be a global variable in" + "%<#pragma acc declare link%>", + decl); + error = true; + continue; + } + break; + + default: + if (global_bindings_p ()) + { + error_at (loc, "invalid OpenACC clause at file scope"); + error = true; + continue; + } + if (DECL_EXTERNAL (decl)) + { + error_at (loc, + "invalid use of %<extern%> variable %qD " + "in %<#pragma acc declare%>", decl); + error = true; + continue; + } + else if (TREE_PUBLIC (decl)) + { + error_at (loc, + "invalid use of %<global%> variable %qD " + "in %<#pragma acc declare%>", decl); + error = true; + continue; + } + break; + } + + if (lookup_attribute ("omp declare target", DECL_ATTRIBUTES (decl)) + || lookup_attribute ("omp declare target link", + DECL_ATTRIBUTES (decl))) + { + error_at (loc, "variable %qD used more than once with " + "%<#pragma acc declare%>", decl); + error = true; + continue; + } + + if (!error) + { + tree id; + + if (OMP_CLAUSE_MAP_KIND (t) == GOMP_MAP_LINK) + id = get_identifier ("omp declare target link"); + else + id = get_identifier ("omp declare target"); + + DECL_ATTRIBUTES (decl) + = tree_cons (id, NULL_TREE, DECL_ATTRIBUTES (decl)); + + if (global_bindings_p ()) + { + symtab_node *node = symtab_node::get (decl); + if (node != NULL) + { + node->offloadable = 1; +#ifdef ENABLE_OFFLOADING + g->have_offload = true; + if (is_a <varpool_node *> (node)) + { + vec_safe_push (offload_vars, decl); + node->force_output = 1; + } +#endif + } + } + } + } + + if (error || global_bindings_p ()) + return; + + stmt = make_node (OACC_DECLARE); + TREE_TYPE (stmt) = void_type_node; + OACC_DECLARE_CLAUSES (stmt) = clauses; + SET_EXPR_LOCATION (stmt, pragma_loc); + + add_stmt (stmt); + + return; +} + +/* OpenACC 2.0: # pragma acc enter data oacc-enter-data-clause[optseq] new-line or diff --git a/gcc/cp/ChangeLog b/gcc/cp/ChangeLog index 9a02de2..7bf18f6 100644 --- a/gcc/cp/ChangeLog +++ b/gcc/cp/ChangeLog @@ -1,3 +1,17 @@ +2015-11-12 James Norris <jnorris@codesourcery.com> + Joseph Myers <joseph@codesourcery.com> + + * parser.c (cp_parser_omp_clause_name): Handle 'device_resident' + clause. + (cp_parser_oacc_data_clause): Handle PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT + and PRAGMA_OMP_CLAUSE_LINK. + (cp_paser_oacc_all_clauses): Handle PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT + and PRAGMA_OMP_CLAUSE_LINK. + (OACC_DECLARE_CLAUSE_MASK): New definition. + (cp_parser_oacc_declare): New function. + (cp_parser_pragma): Handle PRAGMA_OACC_DECLARE. + * pt.c (tsubst_expr): Handle OACC_DECLARE. + 2015-11-12 Jason Merrill <jason@redhat.com> * pt.c (check_explicit_specialization): Check the namespace after diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c index a87675e..0ab5275 100644 --- a/gcc/cp/parser.c +++ b/gcc/cp/parser.c @@ -29128,6 +29128,8 @@ cp_parser_omp_clause_name (cp_parser *parser) result = PRAGMA_OMP_CLAUSE_DEVICE; else if (!strcmp ("deviceptr", p)) result = PRAGMA_OACC_CLAUSE_DEVICEPTR; + else if (!strcmp ("device_resident", p)) + result = PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT; else if (!strcmp ("dist_schedule", p)) result = PRAGMA_OMP_CLAUSE_DIST_SCHEDULE; break; @@ -29541,10 +29543,16 @@ cp_parser_oacc_data_clause (cp_parser *parser, pragma_omp_clause c_kind, case PRAGMA_OACC_CLAUSE_DEVICE: kind = GOMP_MAP_FORCE_TO; break; + case PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT: + kind = GOMP_MAP_DEVICE_RESIDENT; + break; case PRAGMA_OACC_CLAUSE_HOST: case PRAGMA_OACC_CLAUSE_SELF: kind = GOMP_MAP_FORCE_FROM; break; + case PRAGMA_OACC_CLAUSE_LINK: + kind = GOMP_MAP_LINK; + break; case PRAGMA_OACC_CLAUSE_PRESENT: kind = GOMP_MAP_FORCE_PRESENT; break; @@ -31545,6 +31553,10 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask, clauses = cp_parser_oacc_data_clause_deviceptr (parser, clauses); c_name = "deviceptr"; break; + case PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT: + clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses); + c_name = "device_resident"; + break; case PRAGMA_OACC_CLAUSE_FIRSTPRIVATE: clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_FIRSTPRIVATE, clauses); @@ -31569,6 +31581,10 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask, clauses, here); c_name = "independent"; break; + case PRAGMA_OACC_CLAUSE_LINK: + clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses); + c_name = "link"; + break; case PRAGMA_OACC_CLAUSE_NUM_GANGS: code = OMP_CLAUSE_NUM_GANGS; c_name = "num_gangs"; @@ -34526,6 +34542,158 @@ cp_parser_oacc_data (cp_parser *parser, cp_token *pragma_tok) } /* OpenACC 2.0: + # pragma acc declare oacc-data-clause[optseq] new-line +*/ + +#define OACC_DECLARE_CLAUSE_MASK \ + ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPY) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_LINK) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE)) + +static tree +cp_parser_oacc_declare (cp_parser *parser, cp_token *pragma_tok) +{ + tree clauses, stmt, t; + bool error = false; + + clauses = cp_parser_oacc_all_clauses (parser, OACC_DECLARE_CLAUSE_MASK, + "#pragma acc declare", pragma_tok, true); + + + if (find_omp_clause (clauses, OMP_CLAUSE_MAP) == NULL_TREE) + { + error_at (pragma_tok->location, + "no valid clauses specified in %<#pragma acc declare%>"); + return NULL_TREE; + } + + for (tree t = clauses; t; t = OMP_CLAUSE_CHAIN (t)) + { + location_t loc = OMP_CLAUSE_LOCATION (t); + tree decl = OMP_CLAUSE_DECL (t); + if (!DECL_P (decl)) + { + error_at (loc, "array section in %<#pragma acc declare%>"); + error = true; + continue; + } + gcc_assert (OMP_CLAUSE_CODE (t) == OMP_CLAUSE_MAP); + switch (OMP_CLAUSE_MAP_KIND (t)) + { + case GOMP_MAP_FORCE_ALLOC: + case GOMP_MAP_FORCE_TO: + case GOMP_MAP_FORCE_DEVICEPTR: + case GOMP_MAP_DEVICE_RESIDENT: + break; + + case GOMP_MAP_POINTER: + /* Generated by c_finish_omp_clauses from array sections; + avoid spurious diagnostics. */ + break; + + case GOMP_MAP_LINK: + if (!global_bindings_p () + && (TREE_STATIC (decl) + || !DECL_EXTERNAL (decl))) + { + error_at (loc, + "%qD must be a global variable in" + "%<#pragma acc declare link%>", + decl); + error = true; + continue; + } + break; + + default: + if (global_bindings_p ()) + { + error_at (loc, "invalid OpenACC clause at file scope"); + error = true; + continue; + } + if (DECL_EXTERNAL (decl)) + { + error_at (loc, + "invalid use of %<extern%> variable %qD " + "in %<#pragma acc declare%>", decl); + error = true; + continue; + } + else if (TREE_PUBLIC (decl)) + { + error_at (loc, + "invalid use of %<global%> variable %qD " + "in %<#pragma acc declare%>", decl); + error = true; + continue; + } + break; + } + + if (lookup_attribute ("omp declare target", DECL_ATTRIBUTES (decl)) + || lookup_attribute ("omp declare target link", + DECL_ATTRIBUTES (decl))) + { + error_at (loc, "variable %qD used more than once with " + "%<#pragma acc declare%>", decl); + error = true; + continue; + } + + if (!error) + { + tree id; + + if (OMP_CLAUSE_MAP_KIND (t) == GOMP_MAP_LINK) + id = get_identifier ("omp declare target link"); + else + id = get_identifier ("omp declare target"); + + DECL_ATTRIBUTES (decl) + = tree_cons (id, NULL_TREE, DECL_ATTRIBUTES (decl)); + if (global_bindings_p ()) + { + symtab_node *node = symtab_node::get (decl); + if (node != NULL) + { + node->offloadable = 1; +#ifdef ENABLE_OFFLOADING + g->have_offload = true; + if (is_a <varpool_node *> (node)) + { + vec_safe_push (offload_vars, decl); + node->force_output = 1; + } +#endif + } + } + } + } + + if (error || global_bindings_p ()) + return NULL_TREE; + + stmt = make_node (OACC_DECLARE); + TREE_TYPE (stmt) = void_type_node; + OACC_DECLARE_CLAUSES (stmt) = clauses; + SET_EXPR_LOCATION (stmt, pragma_tok->location); + + add_stmt (stmt); + + return NULL_TREE; +} + +/* OpenACC 2.0: # pragma acc enter data oacc-enter-data-clause[optseq] new-line or @@ -36354,6 +36522,10 @@ cp_parser_pragma (cp_parser *parser, enum pragma_context context) cp_parser_omp_declare (parser, pragma_tok, context); return false; + case PRAGMA_OACC_DECLARE: + cp_parser_oacc_declare (parser, pragma_tok); + return false; + case PRAGMA_OACC_ROUTINE: cp_parser_oacc_routine (parser, pragma_tok, context); return false; diff --git a/gcc/cp/pt.c b/gcc/cp/pt.c index 2e3d48b..f3b0cd0 100644 --- a/gcc/cp/pt.c +++ b/gcc/cp/pt.c @@ -15408,6 +15408,14 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl, add_stmt (t); break; + case OACC_DECLARE: + t = copy_node (t); + tmp = tsubst_omp_clauses (OACC_DECLARE_CLAUSES (t), false, false, + args, complain, in_decl); + OACC_DECLARE_CLAUSES (t) = tmp; + add_stmt (t); + break; + case OMP_TARGET_UPDATE: case OMP_TARGET_ENTER_DATA: case OMP_TARGET_EXIT_DATA: diff --git a/gcc/gimple-pretty-print.c b/gcc/gimple-pretty-print.c index 7b50cdf..7764201 100644 --- a/gcc/gimple-pretty-print.c +++ b/gcc/gimple-pretty-print.c @@ -1353,6 +1353,9 @@ dump_gimple_omp_target (pretty_printer *buffer, gomp_target *gs, case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA: kind = " oacc_enter_exit_data"; break; + case GF_OMP_TARGET_KIND_OACC_DECLARE: + kind = " oacc_declare"; + break; default: gcc_unreachable (); } diff --git a/gcc/gimple.h b/gcc/gimple.h index 781801b..e45162d 100644 --- a/gcc/gimple.h +++ b/gcc/gimple.h @@ -170,6 +170,7 @@ enum gf_mask { GF_OMP_TARGET_KIND_OACC_DATA = 7, GF_OMP_TARGET_KIND_OACC_UPDATE = 8, GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA = 9, + GF_OMP_TARGET_KIND_OACC_DECLARE = 10, /* True on an GIMPLE_OMP_RETURN statement if the return does not require a thread synchronization via some sort of barrier. The exact barrier @@ -6004,6 +6005,7 @@ is_gimple_omp_oacc (const gimple *stmt) case GF_OMP_TARGET_KIND_OACC_DATA: case GF_OMP_TARGET_KIND_OACC_UPDATE: case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA: + case GF_OMP_TARGET_KIND_OACC_DECLARE: return true; default: return false; diff --git a/gcc/gimplify.c b/gcc/gimplify.c index f5bd637..3c8f8a2 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -176,6 +176,7 @@ static struct gimplify_omp_ctx *gimplify_omp_ctxp; /* Forward declaration. */ static enum gimplify_status gimplify_compound_expr (tree *, gimple_seq *, bool); +static hash_map<tree, tree> *oacc_declare_returns; /* Shorter alias name for the above function for use in gimplify.c only. */ @@ -1078,6 +1079,7 @@ gimplify_bind_expr (tree *expr_p, gimple_seq *pre_p) gimple_seq body, cleanup; gcall *stack_save; location_t start_locus = 0, end_locus = 0; + tree ret_clauses = NULL; tree temp = voidify_wrapper_expr (bind_expr, NULL); @@ -1179,9 +1181,39 @@ gimplify_bind_expr (tree *expr_p, gimple_seq *pre_p) clobber_stmt = gimple_build_assign (t, clobber); gimple_set_location (clobber_stmt, end_locus); gimplify_seq_add_stmt (&cleanup, clobber_stmt); + + if (flag_openacc && oacc_declare_returns != NULL) + { + tree *c = oacc_declare_returns->get (t); + if (c != NULL) + { + if (ret_clauses) + OMP_CLAUSE_CHAIN (*c) = ret_clauses; + + ret_clauses = *c; + + oacc_declare_returns->remove (t); + + if (oacc_declare_returns->elements () == 0) + { + delete oacc_declare_returns; + oacc_declare_returns = NULL; + } + } + } } } + if (ret_clauses) + { + gomp_target *stmt; + gimple_stmt_iterator si = gsi_start (cleanup); + + stmt = gimple_build_omp_target (NULL, GF_OMP_TARGET_KIND_OACC_DECLARE, + ret_clauses); + gsi_insert_seq_before_without_update (&si, stmt, GSI_NEW_STMT); + } + if (cleanup) { gtry *gs; @@ -5809,6 +5841,26 @@ omp_notice_threadprivate_variable (struct gimplify_omp_ctx *ctx, tree decl, return false; } +/* Return true if global var DECL is device resident. */ + +static bool +device_resident_p (tree decl) +{ + tree attr = lookup_attribute ("oacc declare target", DECL_ATTRIBUTES (decl)); + + if (!attr) + return false; + + for (tree t = TREE_VALUE (attr); t; t = TREE_PURPOSE (t)) + { + tree c = TREE_VALUE (t); + if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DEVICE_RESIDENT) + return true; + } + + return false; +} + /* Determine outer default flags for DECL mentioned in an OMP region but not declared in an enclosing clause. @@ -5908,6 +5960,15 @@ static unsigned oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags) { const char *rkind; + bool on_device = false; + + if ((ctx->region_type & (ORT_ACC_PARALLEL | ORT_ACC_KERNELS)) != 0 + && is_global_var (decl) + && device_resident_p (decl)) + { + on_device = true; + flags |= GOVD_MAP_TO_ONLY; + } switch (ctx->region_type) { @@ -5928,7 +5989,7 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags) || POINTER_TYPE_P (type)) type = TREE_TYPE (type); - if (AGGREGATE_TYPE_P (type)) + if (on_device || AGGREGATE_TYPE_P (type)) /* Aggregates default to 'present_or_copy'. */ flags |= GOVD_MAP; else @@ -7822,6 +7883,121 @@ gimplify_oacc_cache (tree *expr_p, gimple_seq *pre_p) *expr_p = NULL_TREE; } +/* Helper function of gimplify_oacc_declare. The helper's purpose is to, + if required, translate 'kind' in CLAUSE into an 'entry' kind and 'exit' + kind. The entry kind will replace the one in CLAUSE, while the exit + kind will be used in a new omp_clause and returned to the caller. */ + +static tree +gimplify_oacc_declare_1 (tree clause) +{ + HOST_WIDE_INT kind, new_op; + bool ret = false; + tree c = NULL; + + kind = OMP_CLAUSE_MAP_KIND (clause); + + switch (kind) + { + case GOMP_MAP_ALLOC: + case GOMP_MAP_FORCE_ALLOC: + case GOMP_MAP_FORCE_TO: + new_op = GOMP_MAP_FORCE_DEALLOC; + ret = true; + break; + + case GOMP_MAP_FORCE_FROM: + OMP_CLAUSE_SET_MAP_KIND (clause, GOMP_MAP_FORCE_ALLOC); + new_op = GOMP_MAP_FORCE_FROM; + ret = true; + break; + + case GOMP_MAP_FORCE_TOFROM: + OMP_CLAUSE_SET_MAP_KIND (clause, GOMP_MAP_FORCE_TO); + new_op = GOMP_MAP_FORCE_FROM; + ret = true; + break; + + case GOMP_MAP_FROM: + OMP_CLAUSE_SET_MAP_KIND (clause, GOMP_MAP_FORCE_ALLOC); + new_op = GOMP_MAP_FROM; + ret = true; + break; + + case GOMP_MAP_TOFROM: + OMP_CLAUSE_SET_MAP_KIND (clause, GOMP_MAP_TO); + new_op = GOMP_MAP_FROM; + ret = true; + break; + + case GOMP_MAP_DEVICE_RESIDENT: + case GOMP_MAP_FORCE_DEVICEPTR: + case GOMP_MAP_FORCE_PRESENT: + case GOMP_MAP_LINK: + case GOMP_MAP_POINTER: + case GOMP_MAP_TO: + break; + + default: + gcc_unreachable (); + break; + } + + if (ret) + { + c = build_omp_clause (OMP_CLAUSE_LOCATION (clause), OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (c, new_op); + OMP_CLAUSE_DECL (c) = OMP_CLAUSE_DECL (clause); + } + + return c; +} + +/* Gimplify OACC_DECLARE. */ + +static void +gimplify_oacc_declare (tree *expr_p, gimple_seq *pre_p) +{ + tree expr = *expr_p; + gomp_target *stmt; + tree clauses, t; + + clauses = OACC_DECLARE_CLAUSES (expr); + + gimplify_scan_omp_clauses (&clauses, pre_p, ORT_TARGET_DATA, OACC_DECLARE); + + for (t = clauses; t; t = OMP_CLAUSE_CHAIN (t)) + { + tree decl = OMP_CLAUSE_DECL (t); + + if (TREE_CODE (decl) == MEM_REF) + continue; + + if (TREE_CODE (decl) == VAR_DECL + && !is_global_var (decl) + && DECL_CONTEXT (decl) == current_function_decl) + { + tree c = gimplify_oacc_declare_1 (t); + if (c) + { + if (oacc_declare_returns == NULL) + oacc_declare_returns = new hash_map<tree, tree>; + + oacc_declare_returns->put (decl, c); + } + } + + omp_add_variable (gimplify_omp_ctxp, decl, GOVD_SEEN); + } + + stmt = gimple_build_omp_target (NULL, GF_OMP_TARGET_KIND_OACC_DECLARE, + clauses); + + gimplify_seq_add_stmt (pre_p, stmt); + + *expr_p = NULL_TREE; +} + /* Gimplify the contents of an OMP_PARALLEL statement. This involves gimplification of the body, as well as scanning the body for used variables. We need to do this scan now, because variable-sized @@ -10182,11 +10358,15 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p, break; case OACC_HOST_DATA: - case OACC_DECLARE: sorry ("directive not yet implemented"); ret = GS_ALL_DONE; break; + case OACC_DECLARE: + gimplify_oacc_declare (expr_p, pre_p); + ret = GS_ALL_DONE; + break; + case OACC_DATA: case OACC_KERNELS: case OACC_PARALLEL: diff --git a/gcc/omp-builtins.def b/gcc/omp-builtins.def index 0b6bd58..d540dab 100644 --- a/gcc/omp-builtins.def +++ b/gcc/omp-builtins.def @@ -353,3 +353,5 @@ DEF_GOMP_BUILTIN (BUILT_IN_GOMP_TARGET_ENTER_EXIT_DATA, BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_UINT_PTR, ATTR_NOTHROW_LIST) DEF_GOMP_BUILTIN (BUILT_IN_GOMP_TEAMS, "GOMP_teams", BT_FN_VOID_UINT_UINT, ATTR_NOTHROW_LIST) +DEF_GOACC_BUILTIN (BUILT_IN_GOACC_DECLARE, "GOACC_declare", + BT_FN_VOID_INT_SIZE_PTR_PTR_PTR, ATTR_NOTHROW_LIST) diff --git a/gcc/omp-low.c b/gcc/omp-low.c index 51b471c..f7584de 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -12454,6 +12454,7 @@ expand_omp_target (struct omp_region *region) case GF_OMP_TARGET_KIND_OACC_KERNELS: case GF_OMP_TARGET_KIND_OACC_UPDATE: case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA: + case GF_OMP_TARGET_KIND_OACC_DECLARE: data_region = false; break; case GF_OMP_TARGET_KIND_DATA: @@ -12697,6 +12698,9 @@ expand_omp_target (struct omp_region *region) case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA: start_ix = BUILT_IN_GOACC_ENTER_EXIT_DATA; break; + case GF_OMP_TARGET_KIND_OACC_DECLARE: + start_ix = BUILT_IN_GOACC_DECLARE; + break; default: gcc_unreachable (); } @@ -12819,6 +12823,7 @@ expand_omp_target (struct omp_region *region) switch (start_ix) { case BUILT_IN_GOACC_DATA_START: + case BUILT_IN_GOACC_DECLARE: case BUILT_IN_GOMP_TARGET_DATA: break; case BUILT_IN_GOMP_TARGET: @@ -13133,6 +13138,7 @@ build_omp_regions_1 (basic_block bb, struct omp_region *parent, case GF_OMP_TARGET_KIND_EXIT_DATA: case GF_OMP_TARGET_KIND_OACC_UPDATE: case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA: + case GF_OMP_TARGET_KIND_OACC_DECLARE: /* ..., other than for those stand-alone directives... */ region = NULL; break; @@ -14916,6 +14922,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) case GF_OMP_TARGET_KIND_OACC_KERNELS: case GF_OMP_TARGET_KIND_OACC_UPDATE: case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA: + case GF_OMP_TARGET_KIND_OACC_DECLARE: data_region = false; break; case GF_OMP_TARGET_KIND_DATA: @@ -14987,6 +14994,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) case GOMP_MAP_FORCE_TOFROM: case GOMP_MAP_FORCE_PRESENT: case GOMP_MAP_FORCE_DEVICEPTR: + case GOMP_MAP_DEVICE_RESIDENT: + case GOMP_MAP_LINK: gcc_assert (is_gimple_omp_oacc (stmt)); break; default: @@ -16713,6 +16722,7 @@ make_gimple_omp_edges (basic_block bb, struct omp_region **region, case GF_OMP_TARGET_KIND_EXIT_DATA: case GF_OMP_TARGET_KIND_OACC_UPDATE: case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA: + case GF_OMP_TARGET_KIND_OACC_DECLARE: cur_region = cur_region->outer; break; default: diff --git a/gcc/testsuite/ChangeLog b/gcc/testsuite/ChangeLog index 83b47ea..86054c4 100644 --- a/gcc/testsuite/ChangeLog +++ b/gcc/testsuite/ChangeLog @@ -1,3 +1,9 @@ +2015-11-12 James Norris <jnorris@codesourcery.com> + Joseph Myers <joseph@codesourcery.com> + + * c-c++-common/goacc/declare-1.c: New test. + * c-c++-common/goacc/declare-2.c: Likewise. + 2015-11-12 Christophe Lyon <christophe.lyon@linaro.org> [ARM] Remove neon-testgen.ml and generated tests. diff --git a/gcc/testsuite/c-c++-common/goacc/declare-1.c b/gcc/testsuite/c-c++-common/goacc/declare-1.c new file mode 100644 index 0000000..b036c63 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/declare-1.c @@ -0,0 +1,83 @@ +/* Test valid uses of declare directive. */ +/* { dg-do compile } */ + +int v0; +#pragma acc declare create(v0) + +int v1; +#pragma acc declare copyin(v1) + +int *v2; +#pragma acc declare deviceptr(v2) + +int v3; +#pragma acc declare device_resident(v3) + +int v4; +#pragma acc declare link(v4) + +int v5, v6, v7, v8; +#pragma acc declare create(v5, v6) copyin(v7, v8) + +void +f (void) +{ + int va0; +#pragma acc declare create(va0) + + int va1; +#pragma acc declare copyin(va1) + + int *va2; +#pragma acc declare deviceptr(va2) + + int va3; +#pragma acc declare device_resident(va3) + + extern int ve0; +#pragma acc declare create(ve0) + + extern int ve1; +#pragma acc declare copyin(ve1) + + extern int *ve2; +#pragma acc declare deviceptr(ve2) + + extern int ve3; +#pragma acc declare device_resident(ve3) + + extern int ve4; +#pragma acc declare link(ve4) + + int va5; +#pragma acc declare copy(va5) + + int va6; +#pragma acc declare copyout(va6) + + int va7; +#pragma acc declare present(va7) + + int va8; +#pragma acc declare present_or_copy(va8) + + int va9; +#pragma acc declare present_or_copyin(va9) + + int va10; +#pragma acc declare present_or_copyout(va10) + + int va11; +#pragma acc declare present_or_create(va11) + + a: + { + int va0; +#pragma acc declare create(va0) + if (v1) + goto a; + else + goto b; + } + b:; +} diff --git a/gcc/testsuite/c-c++-common/goacc/declare-2.c b/gcc/testsuite/c-c++-common/goacc/declare-2.c new file mode 100644 index 0000000..d24cb22 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/declare-2.c @@ -0,0 +1,79 @@ +/* Test invalid uses of declare directive. */ +/* { dg-do compile } */ + +#pragma acc declare /* { dg-error "no valid clauses" } */ + +#pragma acc declare create(undeclared) /* { dg-error "undeclared" } */ +/* { dg-error "no valid clauses" "second error" { target *-*-* } 6 } */ + +int v0[10]; +#pragma acc declare create(v0[1:3]) /* { dg-error "array section" } */ + +int v1; +#pragma acc declare create(v1, v1) /* { dg-error "more than once" } */ + +int v2; +#pragma acc declare create(v2) +#pragma acc declare copyin(v2) /* { dg-error "more than once" } */ + +int v3; +#pragma acc declare copy(v3) /* { dg-error "at file scope" } */ + +int v4; +#pragma acc declare copyout(v4) /* { dg-error "at file scope" } */ + +int v5; +#pragma acc declare present(v5) /* { dg-error "at file scope" } */ + +int v6; +#pragma acc declare present_or_copy(v6) /* { dg-error "at file scope" } */ + +int v7; +#pragma acc declare present_or_copyin(v7) /* { dg-error "at file scope" } */ + +int v8; +#pragma acc declare present_or_copyout(v8) /* { dg-error "at file scope" } */ + +int v9; +#pragma acc declare present_or_create(v9) /* { dg-error "at file scope" } */ + +int va10; +#pragma acc declare create (va10) +#pragma acc declare link (va10) /* { dg-error "more than once" } */ + +int va11; +#pragma acc declare link (va11) +#pragma acc declare link (va11) /* { dg-error "more than once" } */ + +int va12; +#pragma acc declare create (va12) link (va12) /* { dg-error "more than once" } */ + +void +f (void) +{ + int va0; +#pragma acc declare link(va0) /* { dg-error "global variable" } */ + + extern int ve0; +#pragma acc declare copy(ve0) /* { dg-error "invalid use of" } */ + + extern int ve1; +#pragma acc declare copyout(ve1) /* { dg-error "invalid use of" } */ + + extern int ve2; +#pragma acc declare present(ve2) /* { dg-error "invalid use of" } */ + + extern int ve3; +#pragma acc declare present_or_copy(ve3) /* { dg-error "invalid use of" } */ + + extern int ve4; +#pragma acc declare present_or_copyin(ve4) /* { dg-error "invalid use of" } */ + + extern int ve5; +#pragma acc declare present_or_copyout(ve5) /* { dg-error "invalid use of" } */ + + extern int ve6; +#pragma acc declare present_or_create(ve6) /* { dg-error "invalid use of" } */ + +#pragma acc declare present (v9) /* { dg-error "invalid use of" } */ +} diff --git a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c index 3f0a4e6..caec760 100644 --- a/gcc/tree-pretty-print.c +++ b/gcc/tree-pretty-print.c @@ -654,6 +654,12 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, int flags) case GOMP_MAP_ALWAYS_POINTER: pp_string (pp, "always_pointer"); break; + case GOMP_MAP_DEVICE_RESIDENT: + pp_string (pp, "device_resident"); + break; + case GOMP_MAP_LINK: + pp_string (pp, "link"); + break; default: gcc_unreachable (); } |