From: James Norris Date: Thu, 12 Nov 2015 22:20:41 +0000 (+0000) Subject: c-pragma.c (oacc_pragmas): Add entry for declare directive. X-Git-Url: https://git.libre-soc.org/?a=commitdiff_plain;h=6e232ba4246ca324a663ec5ddf0ba4db5cf3fbad;p=gcc.git c-pragma.c (oacc_pragmas): Add entry for declare directive. 2015-11-12 James Norris Joseph Myers gcc/c-family/ * 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. gcc/c/ * 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. gcc/cp/ * 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. gcc/ * 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. gcc/testsuite * c-c++-common/goacc/declare-1.c: New test. * c-c++-common/goacc/declare-2.c: Likewise. include/ * gomp-constants.h (enum gomp_map_kind): Add GOMP_MAP_DEVICE_RESIDENT and GOMP_MAP_LINK. libgomp/ * libgomp.map (GOACC_2.0.1): Export GOACC_declare. * oacc-parallel.c (GOACC_declare): New function. * testsuite/libgomp.oacc-c-c++-common/declare-1.c: New test. * testsuite/libgomp.oacc-c-c++-common/declare-2.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/declare-4.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/declare-5.c: Likewise. * testsuite/libgomp.oacc-c++/declare-1.C: Likewise. Co-Authored-By: Joseph Myers From-SVN: r230275 --- diff --git a/gcc/ChangeLog b/gcc/ChangeLog index cb418384a32..5093ce06aba 100644 --- a/gcc/ChangeLog +++ b/gcc/ChangeLog @@ -1,3 +1,26 @@ +2015-11-12 James Norris + Joseph Myers + + * 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 [ARM] Remove neon-testgen.ml and generated tests. diff --git a/gcc/c-family/ChangeLog b/gcc/c-family/ChangeLog index eb4d5bfbae2..5611403ce56 100644 --- a/gcc/c-family/ChangeLog +++ b/gcc/c-family/ChangeLog @@ -1,3 +1,11 @@ +2015-11-12 James Norris + Joseph Myers + + * 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 PR c/68107 diff --git a/gcc/c-family/c-pragma.c b/gcc/c-family/c-pragma.c index f86ed384422..12c3e75ef18 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 afeceff1edb..999ac6794d3 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 0191b455398..02ae07e8b6a 100644 --- a/gcc/c/ChangeLog +++ b/gcc/c/ChangeLog @@ -1,3 +1,15 @@ +2015-11-12 James Norris + Joseph Myers + + * 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 PR c/67784 diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c index 89498254b78..c01d651b297 100644 --- a/gcc/c/c-parser.c +++ b/gcc/c/c-parser.c @@ -1228,6 +1228,7 @@ static vec *c_parser_expr_list (c_parser *, bool, bool, vec **, location_t *, tree *, vec *, 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"; @@ -13217,6 +13238,161 @@ c_parser_oacc_data (location_t loc, c_parser *parser) return stmt; } +/* 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 % variable %qD " + "in %<#pragma acc declare%>", decl); + error = true; + continue; + } + else if (TREE_PUBLIC (decl)) + { + error_at (loc, + "invalid use of % 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 (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 diff --git a/gcc/cp/ChangeLog b/gcc/cp/ChangeLog index 9a02de26c1b..7bf18f69002 100644 --- a/gcc/cp/ChangeLog +++ b/gcc/cp/ChangeLog @@ -1,3 +1,17 @@ +2015-11-12 James Norris + Joseph Myers + + * 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 * pt.c (check_explicit_specialization): Check the namespace after diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c index a87675eb1b7..0ab5275d2d6 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"; @@ -34525,6 +34541,158 @@ cp_parser_oacc_data (cp_parser *parser, cp_token *pragma_tok) return stmt; } +/* 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 % variable %qD " + "in %<#pragma acc declare%>", decl); + error = true; + continue; + } + else if (TREE_PUBLIC (decl)) + { + error_at (loc, + "invalid use of % 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 (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 @@ -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 2e3d48bd297..f3b0cd059de 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 7b50cdfc7b3..7764201ea98 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 781801b7c4a..e45162d24e0 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 f5bd637e775..3c8f8a202d1 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 *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; + + 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 0b6bd58e1b3..d540dab7969 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 51b471cff5a..f7584deb3ac 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 83b47ea823b..86054c4c81f 100644 --- a/gcc/testsuite/ChangeLog +++ b/gcc/testsuite/ChangeLog @@ -1,3 +1,9 @@ +2015-11-12 James Norris + Joseph Myers + + * c-c++-common/goacc/declare-1.c: New test. + * c-c++-common/goacc/declare-2.c: Likewise. + 2015-11-12 Christophe Lyon [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 00000000000..b036c636166 --- /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 00000000000..d24cb2222da --- /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 3f0a4e6cfe0..caec7605814 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 (); } diff --git a/include/ChangeLog b/include/ChangeLog index 004d355228a..15b77c723f4 100644 --- a/include/ChangeLog +++ b/include/ChangeLog @@ -1,3 +1,9 @@ +2015-11-12 James Norris + Joseph Myers + + * gomp-constants.h (enum gomp_map_kind): Add GOMP_MAP_DEVICE_RESIDENT + and GOMP_MAP_LINK. + 2015-11-09 Alan Modra PR gdb/17133 diff --git a/include/gomp-constants.h b/include/gomp-constants.h index 7671dd75fad..dffd631aff6 100644 --- a/include/gomp-constants.h +++ b/include/gomp-constants.h @@ -72,6 +72,11 @@ enum gomp_map_kind POINTER_SIZE_UNITS. */ GOMP_MAP_FORCE_DEVICEPTR = (GOMP_MAP_FLAG_SPECIAL_1 | 0), /* Do not map, copy bits for firstprivate instead. */ + /* OpenACC device_resident. */ + GOMP_MAP_DEVICE_RESIDENT = (GOMP_MAP_FLAG_SPECIAL_1 | 1), + /* OpenACC link. */ + GOMP_MAP_LINK = (GOMP_MAP_FLAG_SPECIAL_1 | 2), + /* Allocate. */ GOMP_MAP_FIRSTPRIVATE = (GOMP_MAP_FLAG_SPECIAL | 0), /* Similarly, but store the value in the pointer rather than pointed by the pointer. */ diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog index b4f708e5e46..4e0cddb4e28 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,3 +1,14 @@ +2015-11-12 James Norris + Joseph Myers + + * libgomp.map (GOACC_2.0.1): Export GOACC_declare. + * oacc-parallel.c (GOACC_declare): New function. + * testsuite/libgomp.oacc-c-c++-common/declare-1.c: New test. + * testsuite/libgomp.oacc-c-c++-common/declare-2.c: Likewise. + * testsuite/libgomp.oacc-c-c++-common/declare-4.c: Likewise. + * testsuite/libgomp.oacc-c-c++-common/declare-5.c: Likewise. + * testsuite/libgomp.oacc-c++/declare-1.C: Likewise. + 2015-11-12 Nathan Sidwell * testsuite/libgomp.oacc-c-c++-common/default-1.c: New. diff --git a/libgomp/libgomp.map b/libgomp/libgomp.map index 39faba95cbd..d16710f44f4 100644 --- a/libgomp/libgomp.map +++ b/libgomp/libgomp.map @@ -392,6 +392,7 @@ GOACC_2.0 { GOACC_2.0.1 { global: + GOACC_declare; GOACC_parallel_keyed; } GOACC_2.0; diff --git a/libgomp/oacc-parallel.c b/libgomp/oacc-parallel.c index 525846b0098..a80ede40d1d 100644 --- a/libgomp/oacc-parallel.c +++ b/libgomp/oacc-parallel.c @@ -501,3 +501,61 @@ GOACC_get_thread_num (void) { return 0; } + +void +GOACC_declare (int device, size_t mapnum, + void **hostaddrs, size_t *sizes, unsigned short *kinds) +{ + int i; + + for (i = 0; i < mapnum; i++) + { + unsigned char kind = kinds[i] & 0xff; + + if (kind == GOMP_MAP_POINTER || kind == GOMP_MAP_TO_PSET) + continue; + + switch (kind) + { + case GOMP_MAP_FORCE_ALLOC: + case GOMP_MAP_FORCE_DEALLOC: + case GOMP_MAP_FORCE_FROM: + case GOMP_MAP_FORCE_TO: + case GOMP_MAP_POINTER: + GOACC_enter_exit_data (device, 1, &hostaddrs[i], &sizes[i], + &kinds[i], 0, 0); + break; + + case GOMP_MAP_FORCE_DEVICEPTR: + break; + + case GOMP_MAP_ALLOC: + if (!acc_is_present (hostaddrs[i], sizes[i])) + GOACC_enter_exit_data (device, 1, &hostaddrs[i], &sizes[i], + &kinds[i], 0, 0); + break; + + case GOMP_MAP_TO: + GOACC_enter_exit_data (device, 1, &hostaddrs[i], &sizes[i], + &kinds[i], 0, 0); + + break; + + case GOMP_MAP_FROM: + kinds[i] = GOMP_MAP_FORCE_FROM; + GOACC_enter_exit_data (device, 1, &hostaddrs[i], &sizes[i], + &kinds[i], 0, 0); + break; + + case GOMP_MAP_FORCE_PRESENT: + if (!acc_is_present (hostaddrs[i], sizes[i])) + gomp_fatal ("[%p,%ld] is not mapped", hostaddrs[i], + (unsigned long) sizes[i]); + break; + + default: + assert (0); + break; + } + } +} diff --git a/libgomp/testsuite/libgomp.oacc-c++/declare-1.C b/libgomp/testsuite/libgomp.oacc-c++/declare-1.C new file mode 100644 index 00000000000..0286955d0c7 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c++/declare-1.C @@ -0,0 +1,31 @@ +/* { dg-do run { target openacc_nvidia_accel_selected } } */ + +#include + +template +T foo() +{ + T a, b; + #pragma acc declare create (a) + + #pragma acc parallel copyout (b) + { + a = 5; + b = a; + } + + return b; +} + +int +main (void) +{ + int rc; + + rc = foo(); + + if (rc != 5) + abort (); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-1.c new file mode 100644 index 00000000000..c63a68dbab7 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-1.c @@ -0,0 +1,122 @@ +/* { dg-do run { target openacc_nvidia_accel_selected } } */ + +#include +#include +#include + +#define N 8 + +void +subr2 (int *a) +{ + int i; + int f[N]; +#pragma acc declare copyout (f) + +#pragma acc parallel copy (a[0:N]) + { + for (i = 0; i < N; i++) + { + f[i] = a[i]; + a[i] = f[i] + f[i] + f[i]; + } + } +} + +void +subr1 (int *a) +{ + int f[N]; +#pragma acc declare copy (f) + +#pragma acc parallel copy (a[0:N]) + { + int i; + + for (i = 0; i < N; i++) + { + f[i] = a[i]; + a[i] = f[i] + f[i]; + } + } +} + +int b[8]; +#pragma acc declare create (b) + +int d[8] = { 1, 2, 3, 4, 5, 6, 7, 8 }; +#pragma acc declare copyin (d) + +int +main (int argc, char **argv) +{ + int a[N]; + int e[N]; +#pragma acc declare create (e) + int i; + + for (i = 0; i < N; i++) + a[i] = i + 1; + + if (!acc_is_present (&b, sizeof (b))) + abort (); + + if (!acc_is_present (&d, sizeof (d))) + abort (); + + if (!acc_is_present (&e, sizeof (e))) + abort (); + +#pragma acc parallel copyin (a[0:N]) + { + for (i = 0; i < N; i++) + { + b[i] = a[i]; + a[i] = b[i]; + } + } + + for (i = 0; i < N; i++) + { + if (a[i] != i + 1) + abort (); + } + +#pragma acc parallel copy (a[0:N]) + { + for (i = 0; i < N; i++) + { + e[i] = a[i] + d[i]; + a[i] = e[i]; + } + } + + for (i = 0; i < N; i++) + { + if (a[i] != (i + 1) * 2) + abort (); + } + + for (i = 0; i < N; i++) + { + a[i] = 1234; + } + + subr1 (&a[0]); + + for (i = 0; i < N; i++) + { + if (a[i] != 1234 * 2) + abort (); + } + + subr2 (&a[0]); + + for (i = 0; i < 1; i++) + { + if (a[i] != 1234 * 6) + abort (); + } + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-2.c new file mode 100644 index 00000000000..2078a33afa9 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-2.c @@ -0,0 +1,64 @@ +/* { dg-do run { target openacc_nvidia_accel_selected } } */ + +#include + +#define N 16 + +float c[N]; +#pragma acc declare device_resident (c) + +#pragma acc routine +float +subr2 (float a) +{ + int i; + + for (i = 0; i < N; i++) + c[i] = 2.0; + + for (i = 0; i < N; i++) + a += c[i]; + + return a; +} + +float b[N]; +#pragma acc declare copyin (b) + +#pragma acc routine +float +subr1 (float a) +{ + int i; + + for (i = 0; i < N; i++) + a += b[i]; + + return a; +} + +int +main (int argc, char **argv) +{ + float a; + int i; + + for (i = 0; i < 16; i++) + b[i] = 1.0; + + a = 0.0; + + a = subr1 (a); + + if (a != 16.0) + abort (); + + a = 0.0; + + a = subr2 (a); + + if (a != 32.0) + abort (); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-4.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-4.c new file mode 100644 index 00000000000..013310ecb7d --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-4.c @@ -0,0 +1,41 @@ +/* { dg-do run { target openacc_nvidia_accel_selected } } */ + +#include +#include + +float b; +#pragma acc declare link (b) + +#pragma acc routine +int +func (int a) +{ + b = a + 1; + + return b; +} + +int +main (int argc, char **argv) +{ + float a; + + a = 2.0; + +#pragma acc parallel copy (a) + { + b = a; + a = 1.0; + a = a + b; + } + + if (a != 3.0) + abort (); + + a = func (a); + + if (a != 4.0) + abort (); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-5.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-5.c new file mode 100644 index 00000000000..38c5de063d9 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-5.c @@ -0,0 +1,15 @@ +/* { dg-do run { target openacc_nvidia_accel_selected } } */ + +#include + +int +main (int argc, char **argv) +{ + int a[8] __attribute__((unused)); + + fprintf (stderr, "CheCKpOInT\n"); +#pragma acc declare present (a) +} + +/* { dg-output "CheCKpOInT" } */ +/* { dg-shouldfail "" } */