On Thu, Nov 19, 2015 at 16:31:15 +0100, Jakub Jelinek wrote:
> On Mon, Nov 16, 2015 at 06:40:43PM +0300, Ilya Verbin wrote:
> > @@ -2009,7 +2010,8 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
> > decl = OMP_CLAUSE_DECL (c);
> > /* Global variables with "omp declare target" attribute
> > don't need to be copied, the receiver side will use them
> > - directly. */
> > + directly. However, global variables with "omp declare target link"
> > + attribute need to be copied. */
> > if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
> > && DECL_P (decl)
> > && ((OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_POINTER
> > @@ -2017,7 +2019,9 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
> > != GOMP_MAP_FIRSTPRIVATE_REFERENCE))
> > || TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE)
> > && is_global_var (maybe_lookup_decl_in_outer_ctx (decl, ctx))
> > - && varpool_node::get_create (decl)->offloadable)
> > + && varpool_node::get_create (decl)->offloadable
> > + && !lookup_attribute ("omp declare target link",
> > + DECL_ATTRIBUTES (decl)))
>
> I wonder if Honza/Richi wouldn't prefer to have this info also
> in cgraph, instead of looking up the attribute in each case.
So should I add a new flag into cgraph?
Also it is used in gimplify_adjust_omp_clauses.
> > + if (var.link_ptr_decl == NULL_TREE)
> > + addr = build_fold_addr_expr (var.decl);
> > + else
> > + {
> > + /* For "omp declare target link" var use address of the pointer
> > + instead of address of the var. */
> > + addr = build_fold_addr_expr (var.link_ptr_decl);
> > + /* Most significant bit of the size marks such vars. */
> > + unsigned HOST_WIDE_INT isize = tree_to_uhwi (size);
> > + isize |= 1ULL << (int_size_in_bytes (const_ptr_type_node) * 8 - 1);
> > + size = wide_int_to_tree (const_ptr_type_node, isize);
> > +
> > + /* FIXME: Remove varpool node of var? */
>
> There is varpool_node::remove (), but not sure if at this point all the
> references are already gone.
Actually removing varpool node here will not remove var from the target code, so
I've added a check in cgraphunit.c before assemble_decl ().
> > +class pass_omp_target_link : public gimple_opt_pass
> > +{
> > +public:
> > + pass_omp_target_link (gcc::context *ctxt)
> > + : gimple_opt_pass (pass_data_omp_target_link, ctxt)
> > + {}
> > +
> > + /* opt_pass methods: */
> > + virtual bool gate (function *fun)
> > + {
> > +#ifdef ACCEL_COMPILER
> > + /* FIXME: Replace globals in target regions too or not? */
> > + return lookup_attribute ("omp declare target",
> > + DECL_ATTRIBUTES (fun->decl));
>
> Certainly in "omp declare target entrypoint" regions too.
Done.
> > +unsigned
> > +pass_omp_target_link::execute (function *fun)
> > +{
> > + basic_block bb;
> > + FOR_EACH_BB_FN (bb, fun)
> > + {
> > + gimple_stmt_iterator gsi;
> > + for (gsi = gsi_start_bb (bb); !gsi_end_p (gsi); gsi_next (&gsi))
> > + {
> > + unsigned i;
> > + gimple *stmt = gsi_stmt (gsi);
> > + for (i = 0; i < gimple_num_ops (stmt); i++)
> > + {
> > + tree op = gimple_op (stmt, i);
> > + tree var = NULL_TREE;
> > +
> > + if (!op)
> > + continue;
> > + if (TREE_CODE (op) == VAR_DECL)
> > + var = op;
> > + else if (TREE_CODE (op) == ADDR_EXPR)
> > + {
> > + tree op1 = TREE_OPERAND (op, 0);
> > + if (TREE_CODE (op1) == VAR_DECL)
> > + var = op1;
> > + }
> > + /* FIXME: Support arrays. What else? */
>
> We need to support all the references to the variables.
> So, I think this approach is not right.
>
> > +
> > + if (var && lookup_attribute ("omp declare target link",
> > + DECL_ATTRIBUTES (var)))
> > + {
> > + tree type = TREE_TYPE (var);
> > + tree ptype = build_pointer_type (type);
> > +
> > + /* Find var in offload table. */
> > + omp_offload_var *table_entry = NULL;
> > + for (unsigned j = 0; j < vec_safe_length (offload_vars); j++)
> > + if ((*offload_vars)[j].decl == var)
> > + {
> > + table_entry = &(*offload_vars)[j];
> > + break;
> > + }
>
> Plus this would be terribly expensive if there are many variables in
> offload_vars.
> So, what I think should be done instead is that you first somewhere, perhaps
> when streaming in the decls from LTO in ACCEL_COMPILER or so, create
> the artificial link ptr variables for the "omp declare target link"
> global vars and
> SET_DECL_VALUE_EXPR (var, build_simple_mem_ref (link_ptr_var));
> DECL_HAS_VALUE_EXPR_P (var) = 1;
> and then in this pass just walk_gimple_stmt each stmt, with a
> callback that would check for VAR_DECLs with DECL_HAS_VALUE_EXPR_P set
> and in that case check if they are "omp declare target link", and if found
> signal to the caller that the stmt needs to be regimplified, then just
> gimple_regimplify_operands those stmts.
Cool, it works :) However I had to disable 2 checks in
varpool_node::assemble_decl for ACCEL_COMPILER.
> > + gcc_assert (table_entry);
> > +
> > + /* Get or create artificial pointer for the var. */
> > + tree ptr_decl;
> > + if (table_entry->link_ptr_decl != NULL_TREE)
> > + ptr_decl = table_entry->link_ptr_decl;
> > + else
> > + {
> > + /* FIXME: Create a new node instead of copying?
> > + Which info to preserve? */
> > + ptr_decl = copy_node (var);
>
> I think you want a new node instead of copying. You don't really want to
> copy anything, perhaps TREE_USED, and set DECL_NAME to something derived
> from the original name. Make the ptr DECL_ARTIFICIAL and perhaps
> DECL_NAMELESS.
Done.
> > diff --git a/libgomp/target.c b/libgomp/target.c
> > index ef22329..195be43 100644
> > --- a/libgomp/target.c
> > +++ b/libgomp/target.c
> > @@ -78,6 +78,17 @@ static int num_devices;
> > /* Number of GOMP_OFFLOAD_CAP_OPENMP_400 devices. */
> > static int num_devices_openmp;
> >
> > +/* FIXME: Quick and dirty prototype of keeping correspondence between host
> > + address of the object and target address of the artificial link pointer.
> > + Move it to gomp_device_descr, or where? */
> > +struct link_struct
> > +{
> > + uintptr_t host_start;
> > + uintptr_t tgt_link_ptr;
> > +};
> > +static struct link_struct links[100];
> > +static int link_num;
>
> As for the representation, I think one possibility would be to say define
> REFCOUNT_LINK (~(uintptr_t) 1)
> and register at gomp_load_image_to_device time the link vars with that
> refcount instead of REFCOUNT_INFINITY. If k->refcount == REFCOUNT_LINK
> then k->tgt_offset would be the pointer to the artificial pointer variable
> instead of actual mapping; for say pointer lookup purposes
> k->refcount == REFCOUNT_LINK would be treated as not mapped, and
> gomp_map_vars if mapping something over that would simply temporarily
> replace (remove the old splay tree key, add the new one) the REFCOUNT_LINK
> entry
> with the new mapping (and store the pointer). Then for the even when the
> new mapping's refcount drops to zero we need to ensure that we readd the
> REFCOUNT_LINK entry. For that we need to store the old splay_tree_key
> somewhere. Either we can add it to splay_tree_key_s, but then it will be
> around unconditionally for all entries, and splay_tree_node right now is
> nicely power of 2-ish - 8 pointers. Or stick it somewhere in
> struct target_mem_desc, say splay_tree_key *link; and if the tgt has
> tgt->array
> allocated and any of the mappings were previously REFCOUNT_LINK, then you
> could
> either allocate that link array with not_found_cnt elements, or allocate
> together with tgt->array and just point it after the last entry in
> tgt->array. tgt->link[i] would be NULL if tgt->array[i] splay_tree_node_s
> did not replace REFCOUNT_LINK when created, and the old REFCOUNT_LINK entry
> otherwise. When do_unmap or exit_data, before splay_tree_remove you'd
> find corresponding link entry (k should point to &k->tgt->array[X].key
> for some X, so (splay_tree_node) k - k->tgt->array should be X and thus
> splay_tree_key linkk = NULL;
> if (k->tgt->link)
> linkk = k->tgt->link[(splay_tree_node) k - k->tgt->array];
> before
> splay_tree_remove (&devicep->mem_map, k);
> should hopefully give you the splay_tree_key to insert again.
I implemented the first approach, because the second seems more complicated.
Or should I implement the second?
make check-target-libgomp passed, bootstrap in progress. Is it OK?
gcc/c-family/
* c-common.c (c_common_attribute_table): Handle "omp declare target
link" attribute.
gcc/
* cgraphunit.c (output_in_order): Do not assemble "omp declare target
link" variables in ACCEL_COMPILER.
* gimplify.c (gimplify_adjust_omp_clauses): Do not remove mapping of
"omp declare target link" variables.
* lto/lto.c: Include stringpool.h and fold-const.h.
(offload_handle_link_vars): New static function.
(lto_main): Call offload_handle_link_vars.
* omp-low.c (scan_sharing_clauses): Do not remove mapping of "omp
declare target link" variables.
(add_decls_addresses_to_decl_constructor): For "omp declare target link"
variables output address of the artificial pointer instead of address of
the variable. Set most significant bit of the size to mark them.
(pass_data_omp_target_link): New pass_data.
(pass_omp_target_link): New class.
(find_link_var_op): New static function.
(make_pass_omp_target_link): New function.
* passes.def: Add pass_omp_target_link.
* tree-pass.h (make_pass_omp_target_link): Declare.
* varpool.c (varpool_node::assemble_decl): Allow decls with VALUE_EXPR
in ACCEL_COMPILER.
libgomp/
* libgomp.h (REFCOUNT_LINK): Define.
(struct splay_tree_key_s): Add link_key.
* target.c (gomp_map_vars): Treat REFCOUNT_LINK objects as not mapped.
Replace target address of the pointer with target address of newly
mapped object in the splay tree. Set link pointer on target to the
device address of the mapped object.
(gomp_unmap_vars): Restore target address of the pointer in the splay
tree for REFCOUNT_LINK objects after unmapping.
(gomp_load_image_to_device): Set refcount to REFCOUNT_LINK for "omp
declare target link" objects.
(gomp_exit_data): Restore target address of the pointer in the splay
tree for REFCOUNT_LINK objects after unmapping.
* testsuite/libgomp.c/target-link-1.c: New file.
diff --git a/gcc/c-family/c-common.c b/gcc/c-family/c-common.c
index fe0a235..81defd6 100644
--- a/gcc/c-family/c-common.c
+++ b/gcc/c-family/c-common.c
@@ -822,6 +822,8 @@ const struct attribute_spec c_common_attribute_table[] =
handle_simd_attribute, false },
{ "omp declare target", 0, 0, true, false, false,
handle_omp_declare_target_attribute, false },
+ { "omp declare target link", 0, 0, true, false, false,
+ handle_omp_declare_target_attribute, false },
{ "alloc_align", 1, 1, false, true, true,
handle_alloc_align_attribute, false },
{ "assume_aligned", 1, 2, false, true, true,
diff --git a/gcc/cgraphunit.c b/gcc/cgraphunit.c
index f73d9a7..8bc70f0 100644
--- a/gcc/cgraphunit.c
+++ b/gcc/cgraphunit.c
@@ -2204,6 +2204,13 @@ output_in_order (bool no_reorder)
break;
case ORDER_VAR:
+#ifdef ACCEL_COMPILER
+ /* Do not assemble "omp declare target link" vars. */
+ if (DECL_HAS_VALUE_EXPR_P (nodes[i].u.v->decl)
+ && lookup_attribute ("omp declare target link",
+ DECL_ATTRIBUTES (nodes[i].u.v->decl)))
+ break;
+#endif
nodes[i].u.v->assemble_decl ();
break;
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index a3ed378..5a381da 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -7700,7 +7700,9 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, tree
*list_p,
n = splay_tree_lookup (ctx->variables, (splay_tree_key) decl);
if ((ctx->region_type & ORT_TARGET) != 0
&& !(n->value & GOVD_SEEN)
- && GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c)) == 0)
+ && GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c)) == 0
+ && !lookup_attribute ("omp declare target link",
+ DECL_ATTRIBUTES (decl)))
{
remove = true;
/* For struct element mapping, if struct is never referenced
diff --git a/gcc/lto/lto.c b/gcc/lto/lto.c
index 2661491..58f8a68 100644
--- a/gcc/lto/lto.c
+++ b/gcc/lto/lto.c
@@ -49,6 +49,8 @@ along with GCC; see the file COPYING3. If not see
#include "params.h"
#include "ipa-utils.h"
#include "gomp-constants.h"
+#include "stringpool.h"
+#include "fold-const.h"
/* Number of parallel tasks to run, -1 if we want to use GNU Make jobserver.
*/
@@ -3223,6 +3225,37 @@ lto_init (void)
#endif
}
+/* Create artificial pointers for "omp declare target link" vars. */
+
+static void
+offload_handle_link_vars (void)
+{
+#ifdef ACCEL_COMPILER
+ varpool_node *var;
+ FOR_EACH_VARIABLE (var)
+ if (lookup_attribute ("omp declare target link",
+ DECL_ATTRIBUTES (var->decl)))
+ {
+ tree type = build_pointer_type (TREE_TYPE (var->decl));
+ tree link_ptr_var = make_node (VAR_DECL);
+ TREE_TYPE (link_ptr_var) = type;
+ TREE_USED (link_ptr_var) = 1;
+ TREE_STATIC (link_ptr_var) = 1;
+ DECL_MODE (link_ptr_var) = TYPE_MODE (type);
+ DECL_SIZE (link_ptr_var) = TYPE_SIZE (type);
+ DECL_SIZE_UNIT (link_ptr_var) = TYPE_SIZE_UNIT (type);
+ DECL_ARTIFICIAL (link_ptr_var) = 1;
+ tree var_name = DECL_ASSEMBLER_NAME (var->decl);
+ char *new_name
+ = ACONCAT ((IDENTIFIER_POINTER (var_name), "_linkptr", NULL));
+ DECL_NAME (link_ptr_var) = get_identifier (new_name);
+ SET_DECL_ASSEMBLER_NAME (link_ptr_var, DECL_NAME (link_ptr_var));
+ SET_DECL_VALUE_EXPR (var->decl, build_simple_mem_ref (link_ptr_var));
+ DECL_HAS_VALUE_EXPR_P (var->decl) = 1;
+ }
+#endif
+}
+
/* Main entry point for the GIMPLE front end. This front end has
three main personalities:
@@ -3271,6 +3304,8 @@ lto_main (void)
if (!seen_error ())
{
+ offload_handle_link_vars ();
+
/* If WPA is enabled analyze the whole call graph and create an
optimization plan. Otherwise, read in all the function
bodies and continue with optimization. */
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 0d4c6e5..423b2d1 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -2006,7 +2006,8 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
decl = OMP_CLAUSE_DECL (c);
/* Global variables with "omp declare target" attribute
don't need to be copied, the receiver side will use them
- directly. */
+ directly. However, global variables with "omp declare target link"
+ attribute need to be copied. */
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& DECL_P (decl)
&& ((OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_POINTER
@@ -2014,7 +2015,9 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
!= GOMP_MAP_FIRSTPRIVATE_REFERENCE))
|| TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE)
&& is_global_var (maybe_lookup_decl_in_outer_ctx (decl, ctx))
- && varpool_node::get_create (decl)->offloadable)
+ && varpool_node::get_create (decl)->offloadable
+ && !lookup_attribute ("omp declare target link",
+ DECL_ATTRIBUTES (decl)))
break;
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER)
@@ -18480,13 +18483,35 @@ add_decls_addresses_to_decl_constructor (vec<tree,
va_gc> *v_decls,
for (unsigned i = 0; i < len; i++)
{
tree it = (*v_decls)[i];
- bool is_function = TREE_CODE (it) != VAR_DECL;
+ bool is_var = TREE_CODE (it) == VAR_DECL;
+ bool is_link_var
+ = is_var && DECL_HAS_VALUE_EXPR_P (it)
+ && lookup_attribute ("omp declare target link", DECL_ATTRIBUTES (it));
+
+ tree size = NULL_TREE;
+ if (is_var)
+ size = fold_convert (const_ptr_type_node, DECL_SIZE_UNIT (it));
+
+ tree addr;
+ if (!is_link_var)
+ addr = build_fold_addr_expr (it);
+ else
+ {
+ tree value_expr = DECL_VALUE_EXPR (it);
+ tree link_ptr_decl = TREE_OPERAND (value_expr, 0);
+ varpool_node::finalize_decl (link_ptr_decl);
+ /* For "omp declare target link" var use address of the pointer
+ instead of address of the var. */
+ addr = build_fold_addr_expr (link_ptr_decl);
+ /* Most significant bit of the size marks such vars. */
+ unsigned HOST_WIDE_INT isize = tree_to_uhwi (size);
+ isize |= 1ULL << (int_size_in_bytes (const_ptr_type_node) * 8 - 1);
+ size = wide_int_to_tree (const_ptr_type_node, isize);
+ }
- CONSTRUCTOR_APPEND_ELT (v_ctor, NULL_TREE, build_fold_addr_expr (it));
- if (!is_function)
- CONSTRUCTOR_APPEND_ELT (v_ctor, NULL_TREE,
- fold_convert (const_ptr_type_node,
- DECL_SIZE_UNIT (it)));
+ CONSTRUCTOR_APPEND_ELT (v_ctor, NULL_TREE, addr);
+ if (is_var)
+ CONSTRUCTOR_APPEND_ELT (v_ctor, NULL_TREE, size);
}
}
@@ -19723,4 +19748,84 @@ make_pass_oacc_device_lower (gcc::context *ctxt)
return new pass_oacc_device_lower (ctxt);
}
+/* "omp declare target link" handling pass. */
+
+namespace {
+
+const pass_data pass_data_omp_target_link =
+{
+ GIMPLE_PASS, /* type */
+ "omptargetlink", /* name */
+ OPTGROUP_NONE, /* optinfo_flags */
+ TV_NONE, /* tv_id */
+ PROP_ssa, /* properties_required */
+ 0, /* properties_provided */
+ 0, /* properties_destroyed */
+ 0, /* todo_flags_start */
+ TODO_update_ssa, /* todo_flags_finish */
+};
+
+class pass_omp_target_link : public gimple_opt_pass
+{
+public:
+ pass_omp_target_link (gcc::context *ctxt)
+ : gimple_opt_pass (pass_data_omp_target_link, ctxt)
+ {}
+
+ /* opt_pass methods: */
+ virtual bool gate (function *fun)
+ {
+#ifdef ACCEL_COMPILER
+ tree attrs = DECL_ATTRIBUTES (fun->decl);
+ return lookup_attribute ("omp declare target", attrs)
+ || lookup_attribute ("omp target entrypoint", attrs);
+#else
+ (void) fun;
+ return false;
+#endif
+ }
+
+ virtual unsigned execute (function *);
+};
+
+/* Callback for walk_gimple_stmt used to scan for link var operands. */
+
+static tree
+find_link_var_op (tree *tp, int *walk_subtrees, void *)
+{
+ tree t = *tp;
+
+ if (TREE_CODE (t) == VAR_DECL && DECL_HAS_VALUE_EXPR_P (t)
+ && lookup_attribute ("omp declare target link", DECL_ATTRIBUTES (t)))
+ {
+ *walk_subtrees = 0;
+ return t;
+ }
+
+ return NULL_TREE;
+}
+
+unsigned
+pass_omp_target_link::execute (function *fun)
+{
+ basic_block bb;
+ FOR_EACH_BB_FN (bb, fun)
+ {
+ gimple_stmt_iterator gsi;
+ for (gsi = gsi_start_bb (bb); !gsi_end_p (gsi); gsi_next (&gsi))
+ if (walk_gimple_stmt (&gsi, NULL, find_link_var_op, NULL))
+ gimple_regimplify_operands (gsi_stmt (gsi), &gsi);
+ }
+
+ return 0;
+}
+
+} // anon namespace
+
+gimple_opt_pass *
+make_pass_omp_target_link (gcc::context *ctxt)
+{
+ return new pass_omp_target_link (ctxt);
+}
+
#include "gt-omp-low.h"
diff --git a/gcc/passes.def b/gcc/passes.def
index 1702778..46932b2 100644
--- a/gcc/passes.def
+++ b/gcc/passes.def
@@ -153,6 +153,7 @@ along with GCC; see the file COPYING3. If not see
NEXT_PASS (pass_fixup_cfg);
NEXT_PASS (pass_lower_eh_dispatch);
NEXT_PASS (pass_oacc_device_lower);
+ NEXT_PASS (pass_omp_target_link);
NEXT_PASS (pass_all_optimizations);
PUSH_INSERT_PASSES_WITHIN (pass_all_optimizations)
NEXT_PASS (pass_remove_cgraph_callee_edges);
diff --git a/gcc/tree-pass.h b/gcc/tree-pass.h
index dcd2d5e..f6eabe6 100644
--- a/gcc/tree-pass.h
+++ b/gcc/tree-pass.h
@@ -415,6 +415,7 @@ extern gimple_opt_pass *make_pass_lower_omp (gcc::context
*ctxt);
extern gimple_opt_pass *make_pass_diagnose_omp_blocks (gcc::context *ctxt);
extern gimple_opt_pass *make_pass_expand_omp (gcc::context *ctxt);
extern gimple_opt_pass *make_pass_expand_omp_ssa (gcc::context *ctxt);
+extern gimple_opt_pass *make_pass_omp_target_link (gcc::context *ctxt);
extern gimple_opt_pass *make_pass_oacc_device_lower (gcc::context *ctxt);
extern gimple_opt_pass *make_pass_object_sizes (gcc::context *ctxt);
extern gimple_opt_pass *make_pass_strlen (gcc::context *ctxt);
diff --git a/gcc/varpool.c b/gcc/varpool.c
index 36f19a6..cbd1e05 100644
--- a/gcc/varpool.c
+++ b/gcc/varpool.c
@@ -561,17 +561,21 @@ varpool_node::assemble_decl (void)
are not real variables, but just info for debugging and codegen.
Unfortunately at the moment emutls is not updating varpool correctly
after turning real vars into value_expr vars. */
+#ifndef ACCEL_COMPILER
if (DECL_HAS_VALUE_EXPR_P (decl)
&& !targetm.have_tls)
return false;
+#endif
/* Hard register vars do not need to be output. */
if (DECL_HARD_REGISTER (decl))
return false;
+#ifndef ACCEL_COMPILER
gcc_checking_assert (!TREE_ASM_WRITTEN (decl)
&& TREE_CODE (decl) == VAR_DECL
&& !DECL_HAS_VALUE_EXPR_P (decl));
+#endif
if (!in_other_partition
&& !DECL_EXTERNAL (decl))
diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h
index c467f97..ea63248 100644
--- a/libgomp/libgomp.h
+++ b/libgomp/libgomp.h
@@ -817,6 +817,9 @@ struct target_mem_desc {
/* Special value for refcount - infinity. */
#define REFCOUNT_INFINITY (~(uintptr_t) 0)
+/* Special value for refcount - tgt_offset contains target address of the
+ artificial pointer to "omp declare target link" object. */
+#define REFCOUNT_LINK (~(uintptr_t) 1)
struct splay_tree_key_s {
/* Address of the host object. */
@@ -831,6 +834,8 @@ struct splay_tree_key_s {
uintptr_t refcount;
/* Asynchronous reference count. */
uintptr_t async_refcount;
+ /* Pointer to the original mapping of "omp declare target link" object. */
+ splay_tree_key link_key;
};
/* The comparison function. */
diff --git a/libgomp/target.c b/libgomp/target.c
index cf9d0e6..dcbcaaf 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -453,7 +453,7 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t
mapnum,
}
else
n = splay_tree_lookup (mem_map, &cur_node);
- if (n)
+ if (n && n->refcount != REFCOUNT_LINK)
gomp_map_vars_existing (devicep, n, &cur_node, &tgt->list[i],
kind & typemask);
else
@@ -617,11 +617,19 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t
mapnum,
else
k->host_end = k->host_start + sizeof (void *);
splay_tree_key n = splay_tree_lookup (mem_map, k);
- if (n)
+ if (n && n->refcount != REFCOUNT_LINK)
gomp_map_vars_existing (devicep, n, k, &tgt->list[i],
kind & typemask);
else
{
+ k->link_key = NULL;
+ if (n && n->refcount == REFCOUNT_LINK)
+ {
+ /* Replace target address of the pointer with target address
+ of mapped object in the splay tree. */
+ splay_tree_remove (mem_map, n);
+ k->link_key = n;
+ }
size_t align = (size_t) 1 << (kind >> rshift);
tgt->list[i].key = k;
k->tgt = tgt;
@@ -741,6 +749,16 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t
mapnum,
gomp_fatal ("%s: unhandled kind 0x%.2x", __FUNCTION__,
kind);
}
+
+ if (k->link_key)
+ {
+ /* Set link pointer on target to the device address of the
+ mapped object. */
+ void *tgt_addr = (void *) (tgt->tgt_start + k->tgt_offset);
+ devicep->host2dev_func (devicep->target_id,
+ (void *) n->tgt_offset,
+ &tgt_addr, sizeof (void *));
+ }
array++;
}
}
@@ -866,6 +884,9 @@ gomp_unmap_vars (struct target_mem_desc *tgt, bool
do_copyfrom)
if (do_unmap)
{
splay_tree_remove (&devicep->mem_map, k);
+ if (k->link_key)
+ splay_tree_insert (&devicep->mem_map,
+ (splay_tree_node) k->link_key);
if (k->tgt->refcount > 1)
k->tgt->refcount--;
else
@@ -1005,13 +1026,18 @@ gomp_load_image_to_device (struct gomp_device_descr
*devicep, unsigned version,
for (i = 0; i < num_vars; i++)
{
struct addr_pair *target_var = &target_table[num_funcs + i];
- if (target_var->end - target_var->start
- != (uintptr_t) host_var_table[i * 2 + 1])
+ uintptr_t target_size = target_var->end - target_var->start;
+
+ /* Most significant bit of the size marks "omp declare target link"
+ variables. */
+ bool is_link = target_size & (1ULL << (sizeof (uintptr_t) * 8 - 1));
+
+ if (!is_link && (uintptr_t) host_var_table[i * 2 + 1] != target_size)
{
gomp_mutex_unlock (&devicep->lock);
if (is_register_lock)
gomp_mutex_unlock (®ister_lock);
- gomp_fatal ("Can't map target variables (size mismatch)");
+ gomp_fatal ("Cannot map target variables (size mismatch)");
}
splay_tree_key k = &array->key;
@@ -1019,7 +1045,7 @@ gomp_load_image_to_device (struct gomp_device_descr
*devicep, unsigned version,
k->host_end = k->host_start + (uintptr_t) host_var_table[i * 2 + 1];
k->tgt = tgt;
k->tgt_offset = target_var->start;
- k->refcount = REFCOUNT_INFINITY;
+ k->refcount = is_link ? REFCOUNT_LINK : REFCOUNT_INFINITY;
k->async_refcount = 0;
array->left = NULL;
array->right = NULL;
@@ -1632,6 +1658,9 @@ gomp_exit_data (struct gomp_device_descr *devicep, size_t
mapnum,
if (k->refcount == 0)
{
splay_tree_remove (&devicep->mem_map, k);
+ if (k->link_key)
+ splay_tree_insert (&devicep->mem_map,
+ (splay_tree_node) k->link_key);
if (k->tgt->refcount > 1)
k->tgt->refcount--;
else
diff --git a/libgomp/testsuite/libgomp.c/target-link-1.c
b/libgomp/testsuite/libgomp.c/target-link-1.c
new file mode 100644
index 0000000..681677c
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/target-link-1.c
@@ -0,0 +1,63 @@
+struct S { int s, t; };
+
+int a = 1, b = 1;
+double c[27];
+struct S d = { 8888, 8888 };
+#pragma omp declare target link (a) to (b) link (c, d)
+
+int
+foo (void)
+{
+ return a++ + b++;
+}
+
+int
+bar (int n)
+{
+ int *p1 = &a;
+ int *p2 = &b;
+ c[n] += 2.0;
+ d.s -= 2;
+ d.t -= 2;
+ return *p1 + *p2 + d.s + d.t;
+}
+
+#pragma omp declare target (foo, bar)
+
+int
+main ()
+{
+ a = b = 2;
+ d.s = 17;
+ d.t = 18;
+
+ int res, n = 10;
+ #pragma omp target map (to: a, b, c, d) map (from: res)
+ {
+ res = foo () + foo ();
+ c[n] = 3.0;
+ res += bar (n);
+ }
+
+ int shared_mem = 0;
+ #pragma omp target map (alloc: shared_mem)
+ shared_mem = 1;
+
+ if ((shared_mem && res != (2 + 2) + (3 + 3) + (4 + 4 + 15 + 16))
+ || (!shared_mem && res != (2 + 1) + (3 + 2) + (4 + 3 + 15 + 16)))
+ __builtin_abort ();
+
+ #pragma omp target enter data map (to: c)
+ #pragma omp target update from (c)
+ res = (int) (c[n] + 0.5);
+ if ((shared_mem && res != 5) || (!shared_mem && res != 0))
+ __builtin_abort ();
+
+ #pragma omp target map (to: a, b) map (from: res)
+ res = foo ();
+
+ if ((shared_mem && res != 4 + 4) || (!shared_mem && res != 2 + 3))
+ __builtin_abort ();
+
+ return 0;
+}
-- Ilya
