Hi Tobias,
as we discussed off-list, here is the updated patch for the omp declare target
local clause.
This version has the error catching issues you listed in the last review
handled,
and the interaction between 'requires self_maps/unified_shared_memory' adjusted
(basically, the link transformation should not happen for device-local
variables).
All of the tests we discussed should all be included in this patch.
Patch has finished testing on x86_64-linux/nvptx-offloading without regressions.
Okay for mainline?
Chung-Lin
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-1.f90: Adjust test.
* 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 --git a/gcc/c-family/c-pragma.h b/gcc/c-family/c-pragma.h
index 5d28cb45c66..4d1afe93959 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 3c7572b17b9..927aff6d5b8 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 ed234f075ca..4f6e0ccf886 100644
--- a/gcc/c/c-typeck.cc
+++ b/gcc/c/c-typeck.cc
@@ -17792,6 +17792,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))
{
@@ -17986,6 +18002,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 720dec6841e..5b3f41e8bab 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 978838a050a..f9f05eb0923 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 6737e1e3796..d45503d0b38 100644
--- a/gcc/fortran/openmp.cc
+++ b/gcc/fortran/openmp.cc
@@ -6998,26 +6998,29 @@ gfc_match_omp_declare_target (void)
if (n->sym->attr.in_common)
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)
+ else if (n->sym->attr.omp_groupprivate)
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,
- list == OMP_LIST_LINK
- ? "link" : list == OMP_LIST_TO ? "to" : "enter");
+ (list == OMP_LIST_LINK ? "link"
+ : list == OMP_LIST_LOCAL ? "local"
+ : 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 4f68b64fb08..5458237a002 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 93c8cb5e7b0..dc17b658d10 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 c9b4a94255c..f39635712e4 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 901f2d06634..84da463317d 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 00000000000..714eba2cab5
--- /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 a02bb427f4f..f8fb55e75a4 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 00000000000..969ce975a87
--- /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-1.f90
b/gcc/testsuite/gfortran.dg/gomp/groupprivate-1.f90
index 1d9f5492d6a..a2cb2833add 100644
--- a/gcc/testsuite/gfortran.dg/gomp/groupprivate-1.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/groupprivate-1.f90
@@ -9,8 +9,8 @@ module m
! { dg-error "Sorry, OMP GROUPPRIVATE not implemented, used by 'u' declared at
.1." "" { target *-*-* } .-5 }
! { dg-error "Sorry, OMP GROUPPRIVATE not implemented, used by 'k' declared at
.1." "" { target *-*-* } .-6 }
!
- !$omp groupprivate(x, z) device_Type( any )
!$omp declare target local(x) device_type ( any )
+ !$omp groupprivate(x, z) device_Type( any )
!$omp declare target enter( ii) ,local(y), device_type ( host )
!$omp groupprivate(y) device_type( host)
!$omp groupprivate(v) device_type (nohost )
diff --git a/gcc/testsuite/gfortran.dg/gomp/groupprivate-2.f90
b/gcc/testsuite/gfortran.dg/gomp/groupprivate-2.f90
index 2e921ae542f..641fe3a1fe3 100644
--- a/gcc/testsuite/gfortran.dg/gomp/groupprivate-2.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/groupprivate-2.f90
@@ -35,5 +35,5 @@ subroutine sub()
!$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 device_type(host) local(x4) ! { dg-error "List item
'x4' at .1. set in previous OMP GROUPPRIVATE directive to the different
DEVICE_TYPE 'nohost'" }
-
+ ! { dg-error "List item
'x4' at .1. not appear in the 'local' clause as it was previously specified in
a GROUPPRIVATE directive" "" { target *-*-* } .-1 }
end
diff --git a/gcc/tree-core.h b/gcc/tree-core.h
index 1687b1f9467..18330169d83 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 db3713c2f1b..3b7ad2b1818 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 1025535ec5a..b62fba99996 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 00000000000..def822bf654
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/declare-target-local-1.c
@@ -0,0 +1,27 @@
+
+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)
+{
+ int e;
+ #pragma omp target map(from : e)
+ {
+ e = foo (23);
+ }
+
+ #pragma omp target update from(local_var)
+
+ if (e == local_var)
+ __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 00000000000..a770979c6bf
--- /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 00000000000..69d8a549035
--- /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 00000000000..7a40d01c33f
--- /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 00000000000..33f94f398a7
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/declare-target-local-1.f90
@@ -0,0 +1,51 @@
+! { dg-do run }
+! { 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
+
+ !$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
