https://gcc.gnu.org/g:45e402c8ad5cd260d1fe34ae27630625fc8ee098
commit r17-4586-g45e402c8ad5cd260d1fe34ae27630625fc8ee098 Author: Chung-Lin Tang <[email protected]> Date: Tue Sep 22 20:53:47 2026 +0000 OpenMP: Implement local clause for declare target This patch implements the local clause for omp declare target. The main effect is to avoid adding such 'local' variables to offload_vars in the middle/front ends. This effectively creates non-associated versions of the variable on the host/accel device. This is tested by modifying/returning the local-clause variable on the GPU, and testing if "omp target update" does NOT work. Internally, this is represented as an extension of the existing "omp declare target" attribute, with an additional "local" argument into TREE_VALUE(DECL_ATTRIBUTES(var)). This allows value_member() to work as expected when testing if "local" exists, while handling is the same as "omp declare target" in most places. gcc/c-family/ChangeLog: * c-pragma.h (enum pragma_omp_clause): Add PRAGMA_OMP_CLAUSE_LOCAL. gcc/c/ChangeLog: * c-parser.cc (c_parser_omp_clause_name): Add "local" case. (c_parser_omp_all_clauses): Add PRAGMA_OMP_CLAUSE_LOCAL case. (OMP_DECLARE_TARGET_CLAUSE_MASK): Add PRAGMA_OMP_CLAUSE_LOCAL. ( c_parser_omp_declare_target): Handle OMP_CLAUSE_LOCAL. * c-typeck.cc (c_finish_omp_clauses): Handle OMP_CLAUSE_LOCAL. Emit error if local clause variable appears in map clause. gcc/cp/ChangeLog: * parser.cc (cp_parser_omp_clause_name): Add "local" case. (cp_parser_omp_all_clauses): Add PRAGMA_OMP_CLAUSE_LOCAL case. (handle_omp_declare_target_clause): Handle OMP_CLAUSE_LOCAL. (OMP_DECLARE_TARGET_CLAUSE_MASK): Add PRAGMA_OMP_CLAUSE_LOCAL. * semantics.cc (finish_omp_clauses): Handle OMP_CLAUSE_LOCAL. Emit error if local clause variable appears in map clause. gcc/fortran/ChangeLog: * openmp.cc (gfc_match_omp_declare_target): Adjust error handling to handle local clause. * trans-common.cc (build_common_decl): Adjust handling of device_type clause, remove sorry and FIXME of LOCAL clause, handle local clause. * trans-decl.cc (add_attributes_to_decl): Adjust handling of device_type clause, remove sorry of LOCAL clause, adjust creating of "omp declare target [link]" attribute between OpenACC/OpenMP. gcc/ChangeLog: * omp-offload.cc (omp_discover_implicit_declare_target): For self_maps/unified_shared_memory case, adjust to not convert 'omp declare target local' variables to 'link'. * tree-core.h (enum omp_clause_code): Add OMP_CLAUSE_LOCAL. * tree.cc (omp_clause_num_ops): Add entry for OMP_CLAUSE_LOCAL. (omp_clause_code_name): Likewise. * varpool.cc (varpool_node::get_create): Do not add variable to offload_vars if "local" argument is present in "omp declare target". gcc/testsuite/ChangeLog: * c-c++-common/gomp/declare-target-1.c: Add local clause test. * c-c++-common/gomp/declare-target-8.c: New testcase. * gfortran.dg/gomp/declare-target-1.f90: Add local clause test. * gfortran.dg/gomp/declare-target-7.f90: New testcase. * gfortran.dg/gomp/groupprivate-2.f90: Adjust test. libgomp/ChangeLog: * testsuite/libgomp.c-c++-common/declare-target-local-1.c: New test. * testsuite/libgomp.c-c++-common/declare-target-local-2.c: New test. * testsuite/libgomp.c-c++-common/declare-target-local-3.c: New test. * testsuite/libgomp.c-c++-common/declare-target-local-4.c: New test. * testsuite/libgomp.fortran/declare-target-local-1.f90: New test. Diff: --- gcc/c-family/c-pragma.h | 1 + gcc/c/c-parser.cc | 41 ++++++++++++- gcc/c/c-typeck.cc | 17 ++++++ gcc/cp/parser.cc | 40 ++++++++++++- gcc/cp/semantics.cc | 17 ++++++ gcc/fortran/openmp.cc | 18 +++--- gcc/fortran/trans-common.cc | 32 ++++++----- gcc/fortran/trans-decl.cc | 67 +++++++++++++++------- gcc/omp-offload.cc | 9 ++- gcc/testsuite/c-c++-common/gomp/declare-target-1.c | 2 + gcc/testsuite/c-c++-common/gomp/declare-target-8.c | 29 ++++++++++ .../gfortran.dg/gomp/declare-target-1.f90 | 7 ++- .../gfortran.dg/gomp/declare-target-7.f90 | 25 ++++++++ gcc/testsuite/gfortran.dg/gomp/groupprivate-2.f90 | 15 +++-- gcc/tree-core.h | 3 + gcc/tree.cc | 2 + gcc/varpool.cc | 8 ++- .../libgomp.c-c++-common/declare-target-local-1.c | 34 +++++++++++ .../libgomp.c-c++-common/declare-target-local-2.c | 56 ++++++++++++++++++ .../libgomp.c-c++-common/declare-target-local-3.c | 3 + .../libgomp.c-c++-common/declare-target-local-4.c | 47 +++++++++++++++ .../libgomp.fortran/declare-target-local-1.f90 | 55 ++++++++++++++++++ 22 files changed, 467 insertions(+), 61 deletions(-) diff --git a/gcc/c-family/c-pragma.h b/gcc/c-family/c-pragma.h index 5d28cb45c669..4d1afe93959d 100644 --- a/gcc/c-family/c-pragma.h +++ b/gcc/c-family/c-pragma.h @@ -141,6 +141,7 @@ enum pragma_omp_clause { PRAGMA_OMP_CLAUSE_LASTPRIVATE, PRAGMA_OMP_CLAUSE_LINEAR, PRAGMA_OMP_CLAUSE_LINK, + PRAGMA_OMP_CLAUSE_LOCAL, PRAGMA_OMP_CLAUSE_MAP, PRAGMA_OMP_CLAUSE_MERGEABLE, PRAGMA_OMP_CLAUSE_NOCONTEXT, diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc index 3c7572b17b94..927aff6d5b8e 100644 --- a/gcc/c/c-parser.cc +++ b/gcc/c/c-parser.cc @@ -16592,6 +16592,8 @@ c_parser_omp_clause_name (c_parser *parser) result = PRAGMA_OMP_CLAUSE_LINEAR; else if (!strcmp ("link", p)) result = PRAGMA_OMP_CLAUSE_LINK; + else if (!strcmp ("local", p)) + result = PRAGMA_OMP_CLAUSE_LOCAL; break; case 'm': if (!strcmp ("map", p)) @@ -22530,6 +22532,11 @@ c_parser_omp_all_clauses (c_parser *parser, omp_clause_mask mask, = c_parser_omp_var_list_parens (parser, OMP_CLAUSE_LINK, clauses); c_name = "link"; break; + case PRAGMA_OMP_CLAUSE_LOCAL: + clauses + = c_parser_omp_var_list_parens (parser, OMP_CLAUSE_LOCAL, clauses); + c_name = "local"; + break; case PRAGMA_OMP_CLAUSE_TO: if ((mask & (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_LINK)) != 0) { @@ -28877,6 +28884,7 @@ c_maybe_parse_omp_decl (tree decl, tree d) ( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_TO) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_ENTER) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_LINK) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_LOCAL) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICE_TYPE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_INDIRECT)) @@ -28946,6 +28954,11 @@ c_parser_omp_declare_target (c_parser *parser) error_at (OMP_CLAUSE_LOCATION (c), "%qD specified both in declare target %<link%> and %qs" " clauses", t, OMP_CLAUSE_ENTER_TO (c) ? "to" : "enter"); + else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_LOCAL + || value_member (get_identifier ("local"), TREE_VALUE (at2))) + error_at (OMP_CLAUSE_LOCATION (c), + "%qD specified both in declare target %<link%> and " + "%<local%> clauses", t); else error_at (OMP_CLAUSE_LOCATION (c), "%qD specified both in declare target %<link%> and " @@ -28955,6 +28968,11 @@ c_parser_omp_declare_target (c_parser *parser) if (!at1) { DECL_ATTRIBUTES (t) = tree_cons (id, NULL_TREE, DECL_ATTRIBUTES (t)); + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_LOCAL) + TREE_VALUE (DECL_ATTRIBUTES (t)) + = tree_cons (NULL_TREE, get_identifier ("local"), + TREE_VALUE (DECL_ATTRIBUTES (t))); + if (TREE_CODE (t) != FUNCTION_DECL && !is_global_var (t)) continue; @@ -28965,11 +28983,32 @@ c_parser_omp_declare_target (c_parser *parser) if (ENABLE_OFFLOADING) { g->have_offload = true; - if (is_a <varpool_node *> (node)) + if (is_a <varpool_node *> (node) + && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_LOCAL) vec_safe_push (offload_vars, t); } } } + else + { + bool saw_local = value_member (get_identifier ("local"), + TREE_VALUE (at1)); + if (saw_local && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_ENTER) + { + error_at (OMP_CLAUSE_LOCATION (c), + "%qD specified both in declare target %<local%> and " + "%qs clauses", t, + OMP_CLAUSE_ENTER_TO (c) ? "to" : "enter"); + continue; + } + else if (!saw_local && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_LOCAL) + { + error_at (OMP_CLAUSE_LOCATION (c), + "%qD specified both in declare target %<local%> and " + "%<to%> or %<enter%> clauses", t); + continue; + } + } if (TREE_CODE (t) != FUNCTION_DECL) continue; if ((device_type & OMP_CLAUSE_DEVICE_TYPE_HOST) != 0) diff --git a/gcc/c/c-typeck.cc b/gcc/c/c-typeck.cc index ed755ff0e50b..9b01f3daec9b 100644 --- a/gcc/c/c-typeck.cc +++ b/gcc/c/c-typeck.cc @@ -17807,6 +17807,22 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) remove = true; break; } + + tree attr; + if (DECL_P (t) + && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && (attr = lookup_attribute ("omp declare target", + DECL_ATTRIBUTES (t))) + && value_member (get_identifier ("local"), + TREE_VALUE (attr))) + { + error_at (OMP_CLAUSE_LOCATION (c), + "device-local variable %qD cannot appear " + "in map clause", t); + remove = true; + break; + } + /* OpenACC attach / detach clauses must be pointers. */ if (c_oacc_check_attachments (c)) { @@ -18001,6 +18017,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) case OMP_CLAUSE_ENTER: case OMP_CLAUSE_LINK: + case OMP_CLAUSE_LOCAL: t = OMP_CLAUSE_DECL (c); const char *cname; cname = omp_clause_code_name[OMP_CLAUSE_CODE (c)]; diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc index 9c4245bc3f5e..373d3358c25f 100644 --- a/gcc/cp/parser.cc +++ b/gcc/cp/parser.cc @@ -41577,6 +41577,8 @@ cp_parser_omp_clause_name (cp_parser *parser) result = PRAGMA_OMP_CLAUSE_LINEAR; else if (!strcmp ("link", p)) result = PRAGMA_OMP_CLAUSE_LINK; + else if (!strcmp ("local", p)) + result = PRAGMA_OMP_CLAUSE_LOCAL; break; case 'm': if (!strcmp ("map", p)) @@ -47269,6 +47271,10 @@ cp_parser_omp_all_clauses (cp_parser *parser, omp_clause_mask mask, clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_LINK, clauses); c_name = "link"; break; + case PRAGMA_OMP_CLAUSE_LOCAL: + clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_LOCAL, clauses); + c_name = "local"; + break; case PRAGMA_OMP_CLAUSE_TO: if ((mask & (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_LINK)) != 0) { @@ -54608,6 +54614,11 @@ handle_omp_declare_target_clause (tree c, tree t, int device_type, error_at (OMP_CLAUSE_LOCATION (c), "%qD specified both in declare target %<link%> and %qs" " clauses", t, OMP_CLAUSE_ENTER_TO (c) ? "to" : "enter"); + else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_LOCAL + || value_member (get_identifier ("local"), TREE_VALUE (at2))) + error_at (OMP_CLAUSE_LOCATION (c), + "%qD specified both in declare target %<link%> and " + "%<local%> clauses", t); else error_at (OMP_CLAUSE_LOCATION (c), "%qD specified both in declare target %<link%> and " @@ -54617,6 +54628,11 @@ handle_omp_declare_target_clause (tree c, tree t, int device_type, if (!at1) { DECL_ATTRIBUTES (t) = tree_cons (id, NULL_TREE, DECL_ATTRIBUTES (t)); + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_LOCAL) + TREE_VALUE (DECL_ATTRIBUTES (t)) + = tree_cons (NULL_TREE, get_identifier ("local"), + TREE_VALUE (DECL_ATTRIBUTES (t))); + if (TREE_CODE (t) != FUNCTION_DECL && !is_global_var (t)) return true; @@ -54627,11 +54643,32 @@ handle_omp_declare_target_clause (tree c, tree t, int device_type, if (ENABLE_OFFLOADING) { g->have_offload = true; - if (is_a <varpool_node *> (node)) + if (is_a <varpool_node *> (node) + && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_LOCAL) vec_safe_push (offload_vars, t); } } } + else + { + bool saw_local = value_member (get_identifier ("local"), + TREE_VALUE (at1)); + if (saw_local && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_ENTER) + { + error_at (OMP_CLAUSE_LOCATION (c), + "%qD specified both in declare target %<local%> and " + "%qs clauses", t, + OMP_CLAUSE_ENTER_TO (c) ? "to" : "enter"); + return false; + } + else if (!saw_local && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_LOCAL) + { + error_at (OMP_CLAUSE_LOCATION (c), + "%qD specified both in declare target %<local%> and " + "%<to%> or %<enter%> clauses", t); + return false; + } + } if (TREE_CODE (t) != FUNCTION_DECL) return true; if ((device_type & OMP_CLAUSE_DEVICE_TYPE_HOST) != 0) @@ -54682,6 +54719,7 @@ handle_omp_declare_target_clause (tree c, tree t, int device_type, ( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_TO) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_ENTER) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_LINK) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_LOCAL) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICE_TYPE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_INDIRECT)) diff --git a/gcc/cp/semantics.cc b/gcc/cp/semantics.cc index 978838a050a4..f9f05eb09234 100644 --- a/gcc/cp/semantics.cc +++ b/gcc/cp/semantics.cc @@ -9611,6 +9611,22 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) remove = true; break; } + + tree attr; + if (DECL_P (t) + && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && (attr = lookup_attribute ("omp declare target", + DECL_ATTRIBUTES (t))) + && value_member (get_identifier ("local"), + TREE_VALUE (attr))) + { + error_at (OMP_CLAUSE_LOCATION (c), + "device-local variable %qD cannot appear " + "in map clause", t); + remove = true; + break; + } + /* OpenACC attach / detach clauses must be pointers. */ if (cp_oacc_check_attachments (c)) { @@ -9854,6 +9870,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) case OMP_CLAUSE_ENTER: case OMP_CLAUSE_LINK: + case OMP_CLAUSE_LOCAL: t = OMP_CLAUSE_DECL (c); const char *cname; cname = omp_clause_code_name[OMP_CLAUSE_CODE (c)]; diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc index d49fa79b3b4f..a1e89e386fd1 100644 --- a/gcc/fortran/openmp.cc +++ b/gcc/fortran/openmp.cc @@ -6999,25 +6999,27 @@ gfc_match_omp_declare_target (void) gfc_error_now ("OMP DECLARE TARGET variable at %L is an " "element of a COMMON block", &n->where); else if (n->sym->attr.omp_groupprivate && list != OMP_LIST_LOCAL) - gfc_error_now ("List item %qs at %L not appear in the %qs clause " - "as it was previously specified in a GROUPPRIVATE " - "directive", n->sym->name, &n->where, + gfc_error_now ("List item %qs at %L should not appear in the %qs " + "clause, as it was previously specified in a " + "GROUPPRIVATE directive", n->sym->name, &n->where, list == OMP_LIST_LINK ? "link" : list == OMP_LIST_TO ? "to" : "enter"); else if (n->sym->mark) gfc_error_now ("Variable at %L mentioned multiple times in " "clauses of the same OMP DECLARE TARGET directive", &n->where); - else if ((n->sym->attr.omp_declare_target_link - || n->sym->attr.omp_declare_target_local) - && list != OMP_LIST_LINK - && list != OMP_LIST_LOCAL) + else if ((list != OMP_LIST_LINK + && n->sym->attr.omp_declare_target_link) + || (list != OMP_LIST_LOCAL + && n->sym->attr.omp_declare_target_local)) gfc_error_now ("OMP DECLARE TARGET variable at %L previously " "mentioned in %s clause and later in %s clause", &n->where, n->sym->attr.omp_declare_target_link ? "LINK" : "LOCAL", - list == OMP_LIST_TO ? "TO" : "ENTER"); + (list == OMP_LIST_LOCAL ? "LOCAL" + : list == OMP_LIST_LINK ? "LINK" + : list == OMP_LIST_TO ? "TO" : "ENTER")); else if (n->sym->attr.omp_declare_target && (list == OMP_LIST_LINK || list == OMP_LIST_LOCAL)) gfc_error_now ("OMP DECLARE TARGET variable at %L previously " diff --git a/gcc/fortran/trans-common.cc b/gcc/fortran/trans-common.cc index 4f68b64fb08f..5458237a002c 100644 --- a/gcc/fortran/trans-common.cc +++ b/gcc/fortran/trans-common.cc @@ -436,8 +436,6 @@ build_common_decl (gfc_common_head *com, tree union_type, bool is_init) /* If there is no backend_decl for the common block, build it. */ if (decl == NULL_TREE) { - tree omp_clauses = NULL_TREE; - if (com->is_bind_c == 1 && com->binding_label) decl = build_decl (input_location, VAR_DECL, identifier, union_type); else @@ -469,47 +467,51 @@ build_common_decl (gfc_common_head *com, tree union_type, bool is_init) gfc_set_decl_location (decl, &com->where); + tree arg_list = NULL_TREE; if (com->omp_device_type != OMP_DEVICE_TYPE_UNSET) { - tree c = build_omp_clause (UNKNOWN_LOCATION, OMP_CLAUSE_DEVICE_TYPE); + const char *arg_str = NULL; switch (com->omp_device_type) { case OMP_DEVICE_TYPE_HOST: - OMP_CLAUSE_DEVICE_TYPE_KIND (c) = OMP_CLAUSE_DEVICE_TYPE_HOST; + arg_str = "device_type(host)"; break; case OMP_DEVICE_TYPE_NOHOST: - OMP_CLAUSE_DEVICE_TYPE_KIND (c) = OMP_CLAUSE_DEVICE_TYPE_NOHOST; + arg_str = "device_type(nohost)"; break; case OMP_DEVICE_TYPE_ANY: - OMP_CLAUSE_DEVICE_TYPE_KIND (c) = OMP_CLAUSE_DEVICE_TYPE_ANY; + arg_str = "device_type(any)"; break; default: gcc_unreachable (); } - omp_clauses = c; + arg_list = tree_cons (NULL_TREE, get_identifier (arg_str), arg_list); } + /* Also check trans-decl.cc when updating/removing the following; also update f95.c's gfc_gnu_attributes. */ if (com->omp_groupprivate) gfc_error ("Sorry, OMP GROUPPRIVATE not implemented, used by common " "block %</%s/%> declared at %L", com->name, &com->where); - else if (com->omp_declare_target_local) - /* Use 'else if' as groupprivate implies 'local'. */ - gfc_error ("Sorry, OMP DECLARE TARGET with LOCAL clause not implemented" - ", used by common block %</%s/%> declared at %L", - com->name, &com->where); if (com->omp_declare_target_link) DECL_ATTRIBUTES (decl) = tree_cons (get_identifier ("omp declare target link"), - omp_clauses, DECL_ATTRIBUTES (decl)); + arg_list, DECL_ATTRIBUTES (decl)); else if (com->omp_declare_target) DECL_ATTRIBUTES (decl) = tree_cons (get_identifier ("omp declare target"), - omp_clauses, DECL_ATTRIBUTES (decl)); + arg_list, DECL_ATTRIBUTES (decl)); + else if (com->omp_declare_target_local) + { + arg_list = tree_cons (NULL_TREE, get_identifier ("local"), arg_list); + DECL_ATTRIBUTES (decl) + = tree_cons (get_identifier ("omp declare target"), + arg_list, DECL_ATTRIBUTES (decl)); + } if (com->omp_declare_target_link || com->omp_declare_target - /* FIXME: || com->omp_declare_target_local */) + || com->omp_declare_target_local) { /* Add to offload_vars; get_create does so for omp_declare_target and omp_declare_target_local, omp_declare_target_link requires diff --git a/gcc/fortran/trans-decl.cc b/gcc/fortran/trans-decl.cc index 93c8cb5e7b0b..dc17b658d10e 100644 --- a/gcc/fortran/trans-decl.cc +++ b/gcc/fortran/trans-decl.cc @@ -1584,27 +1584,27 @@ add_attributes_to_decl (tree *decl_p, const gfc_symbol *sym) /* FIXME: 'declare_target_link' permits both any and host, but will fail if one sets OMP_CLAUSE_DEVICE_TYPE_KIND. */ + tree arg_list = NULL_TREE; if (sym_attr.omp_device_type != OMP_DEVICE_TYPE_UNSET && !sym_attr.omp_declare_target_link && !sym_attr.omp_declare_target_indirect /* implies 'any' */) { - tree c = build_omp_clause (UNKNOWN_LOCATION, OMP_CLAUSE_DEVICE_TYPE); + const char *arg_str = NULL; switch (sym_attr.omp_device_type) { case OMP_DEVICE_TYPE_HOST: - OMP_CLAUSE_DEVICE_TYPE_KIND (c) = OMP_CLAUSE_DEVICE_TYPE_HOST; + arg_str = "device_type(host)"; break; case OMP_DEVICE_TYPE_NOHOST: - OMP_CLAUSE_DEVICE_TYPE_KIND (c) = OMP_CLAUSE_DEVICE_TYPE_NOHOST; + arg_str = "device_type(nohost)"; break; case OMP_DEVICE_TYPE_ANY: - OMP_CLAUSE_DEVICE_TYPE_KIND (c) = OMP_CLAUSE_DEVICE_TYPE_ANY; + arg_str = "device_type(any)"; break; default: gcc_unreachable (); } - OMP_CLAUSE_CHAIN (c) = clauses; - clauses = c; + arg_list = tree_cons (NULL_TREE, get_identifier (arg_str), arg_list); } /* Also check trans-common.cc when updating/removing the following; @@ -1612,23 +1612,48 @@ add_attributes_to_decl (tree *decl_p, const gfc_symbol *sym) if (sym_attr.omp_groupprivate) gfc_error ("Sorry, OMP GROUPPRIVATE not implemented, " "used by %qs declared at %L", sym->name, &sym->declared_at); - else if (sym_attr.omp_declare_target_local) - /* Use 'else if' as groupprivate implies 'local'. */ - gfc_error ("Sorry, OMP DECLARE TARGET with LOCAL clause not implemented, " - "used by %qs declared at %L", sym->name, &sym->declared_at); bool has_declare = true; - if (sym_attr.omp_declare_target_link - || sym_attr.oacc_declare_link) - list = tree_cons (get_identifier ("omp declare target link"), - clauses, list); - else if (sym_attr.omp_declare_target - || sym_attr.oacc_declare_create - || sym_attr.oacc_declare_copyin - || sym_attr.oacc_declare_deviceptr - || sym_attr.oacc_declare_device_resident) - list = tree_cons (get_identifier ("omp declare target"), - clauses, list); + + if (flag_openmp) + { + if (sym_attr.omp_declare_target_link) + list = tree_cons (get_identifier ("omp declare target link"), + arg_list, list); + else if (sym_attr.omp_declare_target) + list = tree_cons (get_identifier ("omp declare target"), + arg_list, list); + else if (sym_attr.omp_declare_target_local) + { + /* The OpenMP local clause is encoded as an identifier in the + TREE_VALUE of the attribute, itself as a TREE_LIST. This is because + local has very similar handling to normal "enter", just not added + to offload_vars. */ + arg_list = tree_cons (NULL_TREE, get_identifier ("local"), arg_list); + list = tree_cons (get_identifier ("omp declare target"), + arg_list, list); + } + else + has_declare = false; + } + else if (flag_openacc) + { + /* OpenACC has the routine clause list saved in TREE_VALUE of attribute. + This appears consistent with C/C++. Whether OpenMP/OpenACC handling + should be further unified is TBD. */ + if (sym_attr.oacc_declare_link) + list = tree_cons (get_identifier ("omp declare target link"), + clauses, list); + else if (sym_attr.omp_declare_target + || sym_attr.oacc_declare_create + || sym_attr.oacc_declare_copyin + || sym_attr.oacc_declare_deviceptr + || sym_attr.oacc_declare_device_resident) + list = tree_cons (get_identifier ("omp declare target"), + clauses, list); + else + has_declare = false; + } else has_declare = false; diff --git a/gcc/omp-offload.cc b/gcc/omp-offload.cc index c9b4a94255cf..f39635712e48 100644 --- a/gcc/omp-offload.cc +++ b/gcc/omp-offload.cc @@ -399,10 +399,13 @@ omp_discover_implicit_declare_target (void) & (OMP_REQUIRES_SELF_MAPS | OMP_REQUIRES_UNIFIED_SHARED_MEMORY)) FOR_EACH_VARIABLE (vnode) { + tree attr; /* If 'self_maps' or 'unified_shared_memory' is enabled, - remove 'enter/to' and add 'link'. */ - if (lookup_attribute ("omp declare target", - DECL_ATTRIBUTES (vnode->decl))) + remove 'enter/to' and add 'link'. + Note that "declare target local" is preserved and not converted. */ + if ((attr = lookup_attribute ("omp declare target", + DECL_ATTRIBUTES (vnode->decl))) + && !value_member (get_identifier ("local"), TREE_VALUE (attr))) { DECL_ATTRIBUTES (vnode->decl) = remove_attribute ("omp declare target", diff --git a/gcc/testsuite/c-c++-common/gomp/declare-target-1.c b/gcc/testsuite/c-c++-common/gomp/declare-target-1.c index 901f2d06634d..84da463317d1 100644 --- a/gcc/testsuite/c-c++-common/gomp/declare-target-1.c +++ b/gcc/testsuite/c-c++-common/gomp/declare-target-1.c @@ -5,9 +5,11 @@ int foo (void), bar (void); extern int a; int b; char d; +double e; #pragma omp declare target long c; #pragma omp end declare target #pragma omp declare target (bar, a) #pragma omp declare target to (b) link (d) enter (foo) +#pragma omp declare target local (e) diff --git a/gcc/testsuite/c-c++-common/gomp/declare-target-8.c b/gcc/testsuite/c-c++-common/gomp/declare-target-8.c new file mode 100644 index 000000000000..714eba2cab55 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/declare-target-8.c @@ -0,0 +1,29 @@ +/* { dg-do compile } */ +/* { dg-options "-fopenmp" } */ +// { dg-additional-options "-Wno-deprecated-openmp" } + +int var1; +#pragma omp declare target link(var1) +#pragma omp declare target local(var1) /* { dg-error "'var1' specified both in declare target 'link' and 'local' clauses" } */ + +int var2; +#pragma omp declare target local(var2) +#pragma omp declare target link(var2) /* { dg-error "'var2' specified both in declare target 'link' and 'local' clauses" } */ + +int var3; +#pragma omp declare target enter(var3) +#pragma omp declare target local(var3) /* { dg-error "'var3' specified both in declare target 'local' and 'to' or 'enter' clauses" } */ + +int var4; +#pragma omp declare target local(var4) +#pragma omp declare target enter(var4) /* { dg-error "'var4' specified both in declare target 'local' and 'enter' clauses" } */ +#pragma omp declare target to(var4) /* { dg-error "'var4' specified both in declare target 'local' and 'to' clauses" } */ + +#pragma omp begin declare target +int var5; +#pragma omp declare target local(var5) /* { dg-error "'var5' specified both in declare target 'local' and 'to' or 'enter' clauses" } */ +#pragma omp end declare target + +int var6, var7; +#pragma omp declare target local(var6) local(var6) /* { dg-error "'var6' appears more than once on the same 'declare target' directive" } */ +#pragma omp declare target local(var7) enter(var7) /* { dg-error "'var7' appears more than once on the same 'declare target' directive" } */ diff --git a/gcc/testsuite/gfortran.dg/gomp/declare-target-1.f90 b/gcc/testsuite/gfortran.dg/gomp/declare-target-1.f90 index a02bb427f4f6..f8fb55e75a4b 100644 --- a/gcc/testsuite/gfortran.dg/gomp/declare-target-1.f90 +++ b/gcc/testsuite/gfortran.dg/gomp/declare-target-1.f90 @@ -14,14 +14,17 @@ subroutine bar !$omp declare target integer, save :: var_9 !$omp declare target link (var_8) to (baz, var_7) link (var_9) to (var_10) - integer, save :: var_7, var_8, var_10 - integer :: var_11, var_12, var_13, var_14 + integer, save :: var_7, var_8, var_10, var_16 + integer :: var_11, var_12, var_13, var_14, var_15 common /c1/ var_11, var_12 common /c2/ var_13 common /c3/ var_14 + common /c4/ var_15 !$omp declare target (baz, var_7, var_10, /c1/) !$omp declare target to (/c2/) !$omp declare target link (/c3/) !$omp declare target (bar) + !$omp declare target local (var_16) + !$omp declare target local (/c4/) call baz end subroutine diff --git a/gcc/testsuite/gfortran.dg/gomp/declare-target-7.f90 b/gcc/testsuite/gfortran.dg/gomp/declare-target-7.f90 new file mode 100644 index 000000000000..969ce975a876 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/declare-target-7.f90 @@ -0,0 +1,25 @@ +! { dg-do compile } +! { dg-additional-options "-Wno-deprecated-openmp" } + + integer :: var1 + !$omp declare target link(var1) + !$omp declare target local(var1) ! { dg-error "OMP DECLARE TARGET variable at \\(1\\) previously mentioned in LINK clause and later in LOCAL clause" } + + integer :: var2 + !$omp declare target local(var2) + !$omp declare target link(var2) ! { dg-error "OMP DECLARE TARGET variable at \\(1\\) previously mentioned in LOCAL clause and later in LINK clause" } + + integer :: var3 + !$omp declare target enter(var3) + !$omp declare target local(var3) ! { dg-error "OMP DECLARE TARGET variable at \\(1\\) previously mentioned in TO or ENTER clause and later in LOCAL clause" } + + integer :: var4 + !$omp declare target local(var4) + !$omp declare target enter(var4) ! { dg-error "OMP DECLARE TARGET variable at \\(1\\) previously mentioned in LOCAL clause and later in ENTER clause" } + !$omp declare target to(var4) ! { dg-error "OMP DECLARE TARGET variable at \\(1\\) previously mentioned in LOCAL clause and later in TO clause" } + + integer :: var5, var6 + !$omp declare target local(var5) local(var5) ! { dg-error "Variable at \\(1\\) mentioned multiple times in clauses of the same OMP DECLARE TARGET directive" } + !$omp declare target local(var6) enter(var6) ! { dg-error "Variable at \\(1\\) mentioned multiple times in clauses of the same OMP DECLARE TARGET directive" } + +end program diff --git a/gcc/testsuite/gfortran.dg/gomp/groupprivate-2.f90 b/gcc/testsuite/gfortran.dg/gomp/groupprivate-2.f90 index 2e921ae542fa..6ac5ea21af38 100644 --- a/gcc/testsuite/gfortran.dg/gomp/groupprivate-2.f90 +++ b/gcc/testsuite/gfortran.dg/gomp/groupprivate-2.f90 @@ -6,9 +6,9 @@ module m integer :: x, y(20), z, v, q, r,o, b2,c !$omp groupprivate(x, z, o) device_Type( any ) - !$omp declare target enter(x) device_type ( any ) ! { dg-error "List item 'x' at .1. not appear in the 'enter' clause as it was previously specified in a GROUPPRIVATE directive" } - !$omp declare target to(z) device_type ( any ) ! { dg-error "List item 'z' at .1. not appear in the 'to' clause as it was previously specified in a GROUPPRIVATE directive" } - !$omp declare target link(o) device_type ( any ) ! { dg-error "List item 'o' at .1. not appear in the 'link' clause as it was previously specified in a GROUPPRIVATE directive" } + !$omp declare target enter(x) device_type ( any ) ! { dg-error "List item 'x' at .1. should not appear in the 'enter' clause, as it was previously specified in a GROUPPRIVATE directive" } + !$omp declare target to(z) device_type ( any ) ! { dg-error "List item 'z' at .1. should not appear in the 'to' clause, as it was previously specified in a GROUPPRIVATE directive" } + !$omp declare target link(o) device_type ( any ) ! { dg-error "List item 'o' at .1. should not appear in the 'link' clause, as it was previously specified in a GROUPPRIVATE directive" } !$omp declare target enter( ii) ,local(y,c), link(r), to(q) device_type ( host ) !$omp groupprivate(r,q) device_type(host) ! { dg-error "List item 'q' at .1. implies OMP DECLARE TARGET with the LOCAL clause, but it has been specified with a different clause before" "" { target *-*-* } .-1 } @@ -30,10 +30,9 @@ subroutine sub() !$omp groupprivate(x3) device_type (host ) !$omp groupprivate(x4) device_type( nohost) - !$omp declare target(x0) ! { dg-error "List item 'x0' at .1. not appear in the 'enter' clause as it was previously specified in a GROUPPRIVATE directive" } - !$omp declare target device_type(any) to(x1) ! { dg-error "List item 'x1' at .1. not appear in the 'to' clause as it was previously specified in a GROUPPRIVATE directive" } - !$omp declare target device_type(any) enter(x2) ! { dg-error "List item 'x2' at .1. not appear in the 'enter' clause as it was previously specified in a GROUPPRIVATE directive" } - !$omp declare target device_type(host) link(x3) ! { dg-error "List item 'x3' at .1. not appear in the 'link' clause as it was previously specified in a GROUPPRIVATE directive" } + !$omp declare target(x0) ! { dg-error "List item 'x0' at .1. should not appear in the 'enter' clause, as it was previously specified in a GROUPPRIVATE directive" } + !$omp declare target device_type(any) to(x1) ! { dg-error "List item 'x1' at .1. should not appear in the 'to' clause, as it was previously specified in a GROUPPRIVATE directive" } + !$omp declare target device_type(any) enter(x2) ! { dg-error "List item 'x2' at .1. should not appear in the 'enter' clause, as it was previously specified in a GROUPPRIVATE directive" } + !$omp declare target device_type(host) link(x3) ! { dg-error "List item 'x3' at .1. should not appear in the 'link' clause, as it was previously specified in a GROUPPRIVATE directive" } !$omp declare target device_type(host) local(x4) ! { dg-error "List item 'x4' at .1. set in previous OMP GROUPPRIVATE directive to the different DEVICE_TYPE 'nohost'" } - end diff --git a/gcc/tree-core.h b/gcc/tree-core.h index 1687b1f94672..18330169d831 100644 --- a/gcc/tree-core.h +++ b/gcc/tree-core.h @@ -330,6 +330,9 @@ enum omp_clause_code { /* OpenMP clause: link (variable-list). */ OMP_CLAUSE_LINK, + /* OpenMP clause: local (variable-list). */ + OMP_CLAUSE_LOCAL, + /* OpenMP clause: detach (event-handle). */ OMP_CLAUSE_DETACH, diff --git a/gcc/tree.cc b/gcc/tree.cc index db3713c2f1b8..3b7ad2b18180 100644 --- a/gcc/tree.cc +++ b/gcc/tree.cc @@ -318,6 +318,7 @@ unsigned const char omp_clause_num_ops[] = 1, /* OMP_CLAUSE_UNIFORM */ 1, /* OMP_CLAUSE_ENTER */ 1, /* OMP_CLAUSE_LINK */ + 1, /* OMP_CLAUSE_LOCAL */ 1, /* OMP_CLAUSE_DETACH */ 1, /* OMP_CLAUSE_USE_DEVICE_PTR */ 1, /* OMP_CLAUSE_USE_DEVICE_ADDR */ @@ -424,6 +425,7 @@ const char * const omp_clause_code_name[] = "uniform", "enter", "link", + "local", "detach", "use_device_ptr", "use_device_addr", diff --git a/gcc/varpool.cc b/gcc/varpool.cc index 1025535ec5aa..b62fba999968 100644 --- a/gcc/varpool.cc +++ b/gcc/varpool.cc @@ -150,14 +150,18 @@ varpool_node::get_create (tree decl) node = varpool_node::create_empty (); node->decl = decl; + tree attr; if ((flag_openacc || flag_openmp) - && lookup_attribute ("omp declare target", DECL_ATTRIBUTES (decl))) + && (attr = lookup_attribute ("omp declare target", + DECL_ATTRIBUTES (decl)))) { node->offloadable = 1; if (ENABLE_OFFLOADING && !DECL_EXTERNAL (decl)) { g->have_offload = true; - if (!in_lto_p) + if (!in_lto_p + && !value_member (get_identifier ("local"), + TREE_VALUE (attr))) vec_safe_push (offload_vars, decl); } } diff --git a/libgomp/testsuite/libgomp.c-c++-common/declare-target-local-1.c b/libgomp/testsuite/libgomp.c-c++-common/declare-target-local-1.c new file mode 100644 index 000000000000..c6b4d2ecf700 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/declare-target-local-1.c @@ -0,0 +1,34 @@ +/* { dg-do run { target offload_device } } */ + +int local_var = 900; +#pragma omp declare target local(local_var) + +#pragma omp begin declare target +int foo (int x) +{ + local_var += x; + return local_var; +} +#pragma omp end declare target + +int main (void) +{ + /* Variables declared in the 'local' clause should not be corresponding + storage between host/device. Test if 'target update' has no effect + and host-side 'local_var' retains original value. */ + + int e; + #pragma omp target map(from : e) + { + e = foo (23); + } + + #pragma omp target update from(local_var) + + if (local_var == e) + __builtin_abort (); + if (local_var != 900) + __builtin_abort (); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/declare-target-local-2.c b/libgomp/testsuite/libgomp.c-c++-common/declare-target-local-2.c new file mode 100644 index 000000000000..a770979c6bf7 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/declare-target-local-2.c @@ -0,0 +1,56 @@ +#include <assert.h> + +#ifdef USE_SELF_MAPS +#pragma omp self_maps +#endif + +int var[2] = { 5, 6 }; +#pragma omp declare target local(var) + +int +get_var_inc (void) +{ + var[0]++; + return var[0]; +} + +__UINTPTR_TYPE__ +get_addr (void) +{ + return (__UINTPTR_TYPE__) &var[0]; +} + +int main (void) +{ + __UINTPTR_TYPE__ addr_tgt, addr_fn; + bool is_initial_dev; + int val; + var[0] = 7; + + #pragma omp target map(from: val, is_initial_dev, addr_tgt, addr_fn) + { + is_initial_dev = __builtin_omp_is_initial_device (); + addr_tgt = (__UINTPTR_TYPE__) &var[0]; + addr_fn = get_addr (); + + var[0] += 5; + val = get_var_inc (); + } + + if (is_initial_dev) + { + assert (val == var[0]); + assert (val == 13); + assert (&var[0] == (void*) addr_fn); + } + else + { + assert (val == 11); + assert (var[0] == 7); + assert (&var[0] != (void*) addr_fn); + } + + assert (addr_tgt == addr_fn); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/declare-target-local-3.c b/libgomp/testsuite/libgomp.c-c++-common/declare-target-local-3.c new file mode 100644 index 000000000000..69d8a5490358 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/declare-target-local-3.c @@ -0,0 +1,3 @@ + +#define USE_SELF_MAPS +#include "declare-target-local-2.c" diff --git a/libgomp/testsuite/libgomp.c-c++-common/declare-target-local-4.c b/libgomp/testsuite/libgomp.c-c++-common/declare-target-local-4.c new file mode 100644 index 000000000000..7a40d01c33f5 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/declare-target-local-4.c @@ -0,0 +1,47 @@ +#include <omp.h> +#include <assert.h> + +int f (int dev_num) +{ + static int X[1] = { 5 }; + #pragma omp declare target local(X) + + int res; + #pragma omp target map(from: res) device(dev_num) + res = ++X[0]; + + return res; +} + +int main (void) +{ + int cnt; + + cnt = f (omp_initial_device); + assert (cnt == 6); + + cnt = f (omp_initial_device); + assert (cnt == 7); + + if (omp_get_num_devices () > 0) + { + cnt = f (omp_default_device); + assert (cnt == 6); + cnt = f (omp_default_device); + assert (cnt == 7); + cnt = f (omp_default_device); + assert (cnt == 8); + cnt = f (omp_default_device); + assert (cnt == 9); + + cnt = f (omp_initial_device); + assert (cnt == 8); + } + else + { + cnt = f (omp_default_device); + assert (cnt == 8); + } + + return 0; +} diff --git a/libgomp/testsuite/libgomp.fortran/declare-target-local-1.f90 b/libgomp/testsuite/libgomp.fortran/declare-target-local-1.f90 new file mode 100644 index 000000000000..81303db03bad --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/declare-target-local-1.f90 @@ -0,0 +1,55 @@ +! { dg-do run { target offload_device } } +! { dg-additional-options "-Wno-deprecated-openmp" } +module m + implicit none (type, external) + + integer :: local_var = 900 + !$omp declare target local(local_var) + + integer :: local_cblk_var = 800 + common /cblk/ local_cblk_var + !$omp declare target local(/cblk/) + +contains + integer function incr_local_var (x) + !$omp declare target + integer :: x + local_var = local_var + x + incr_local_var = local_var + end function + + integer function incr_local_cblk_var (x) + !$omp declare target + integer :: x + local_cblk_var = local_cblk_var + x + incr_local_cblk_var = local_cblk_var + end function + +end module m + +program main + use m + implicit none (type, external) + integer :: e + + ! Variables declared in the 'local' clause should not be corresponding + ! storage between host/device. Test if 'target update' has no effect + ! and host-side 'local_var' retains original value. */ + + !$omp target map(from : e) + e = incr_local_var (23) + !$omp end target + + !$omp target update from(local_var) + if (local_var == e) stop 1 + if (local_var /= 900) stop 2 + + !$omp target map(from : e) + e = incr_local_cblk_var (68) + !$omp end target + + !$omp target update from(local_cblk_var) + if (local_cblk_var == e) stop 3 + if (local_cblk_var /= 800) stop 4 + +end program main
