[gomp4.5] Handle #pragma omp declare target link
Ilya Verbin
iverbin@gmail.com
Mon Dec 14 17:18:00 GMT 2015
On Mon, Nov 30, 2015 at 21:49:02 +0100, Jakub Jelinek wrote:
> On Mon, Nov 30, 2015 at 11:29:34PM +0300, Ilya Verbin wrote:
> > > This looks wrong, both of these clearly could affect anything with
> > > DECL_HAS_VALUE_EXPR_P, not just the link vars.
> > > So, if you need to handle the "omp declare target link" vars specially,
> > > you should only handle those specially and nothing else. And please try to
> > > explain why.
> >
> > Actually these ifndefs are not needed, because assemble_decl never will be
> > called by accel compiler for original link vars. I've added a check into
> > output_in_order, but missed a second place where assemble_decl is called -
> > symbol_table::output_variables. So, fixed now.
>
> Great.
>
> > > Do we need to do anything in gomp_unload_image_from_device ?
> > > I mean at least in questionable programs that for link vars don't decrement
> > > the refcount of the var that replaced the link var to 0 first before
> > > dlclosing the library.
> > > At least host_var_table[j * 2 + 1] will have the MSB set, so we need to
> > > handle it differently. Perhaps for that case perform a lookup, and if we
> > > get something which has link_map non-NULL, first perform as if there is
> > > target exit data delete (var) on it first?
> >
> > You're right, it doesn't deallocate memory on the device if DSO leaves nonzero
> > refcount. And currently host compiler doesn't set MSB in host_var_table, it's
> > set only by accel compiler. But it's possible to do splay_tree_lookup for each
> > var to determine whether is it linked or not, like in the patch bellow.
> > Or do you prefer to set the bit in host compiler too? It requires
> > lookup_attribute ("omp declare target link") for all vars in the table during
> > compilation, but allows to do splay_tree_lookup at run-time only for vars with
> > MSB set in host_var_table.
> > Unfortunately, calling gomp_exit_data from gomp_unload_image_from_device works
> > only for DSO, but it crashed when an executable leaves nonzero refcount, because
> > target device may be already uninitialized from plugin's __run_exit_handlers
> > (and it is in case of intelmic), so gomp_exit_data cannot run free_func.
> > Is it possible do add some atexit (...) to libgomp, which will set shutting_down
> > flag, and just do nothing in gomp_unload_image_from_device if it is set?
>
> Sorry, I didn't mean you should call gomp_exit_data, what I meant was that
> you perform the same action as would delete(var) do in that case.
> Calling gomp_exit_data e.g. looks it up again etc.
> Supposedly having the MSB in host table too is useful, so if you could
> handle that, it would be nice. And splay_tree_lookup only if the MSB is
> set.
> So,
> if (!host_data_has_msb_set)
> splay_tree_remove (&devicep->mem_map, &k);
> else
> {
> splay_tree_key n = splay_tree_lookup (&devicep->mem_map, &k);
> if (n->link_key)
> {
> n->refcount = 0;
> n->link_key = NULL;
> splay_tree_remove (&devicep->mem_map, n);
> if (n->tgt->refcount > 1)
> n->tgt->refcount--;
> else
> gomp_unmap_tgt (n->tgt);
> }
> else
> splay_tree_remove (&devicep->mem_map, n);
> }
> or so.
Here is an updated patch. Now MSB is set in both tables, and
gomp_unload_image_from_device is changed. I've verified using simple DSO
testcase, that memory on target is freed after dlclose.
bootstrap and make check on x86_64-linux passed.
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 (symbol_table::output_variables): Do not assemble "omp
declare target link" variables 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_unload_image_from_device): Replace j with i. Force unmap of all
"omp declare target link" objects, which were mapped for the image.
(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 9bc02fc..4250cdf 100644
--- a/gcc/c-family/c-common.c
+++ b/gcc/c-family/c-common.c
@@ -821,6 +821,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 3d86c36..8443cb0 100644
--- a/gcc/cgraphunit.c
+++ b/gcc/cgraphunit.c
@@ -2210,6 +2210,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 80c6bf2..438efba 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -7910,7 +7910,9 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, 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 fcf7caf..5fd50dc 100644
--- a/gcc/lto/lto.c
+++ b/gcc/lto/lto.c
@@ -50,6 +50,8 @@ along with GCC; see the file COPYING3. If not see
#include "ipa-utils.h"
#include "gomp-constants.h"
#include "lto-symtab.h"
+#include "stringpool.h"
+#include "fold-const.h"
/* Number of parallel tasks to run, -1 if we want to use GNU Make jobserver. */
@@ -3226,6 +3228,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:
@@ -3274,6 +3307,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 5643480..676b1df 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -2026,7 +2026,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
@@ -2034,7 +2035,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)
@@ -18588,13 +18591,45 @@ 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
+#ifdef ACCEL_COMPILER
+ && DECL_HAS_VALUE_EXPR_P (it)
+#endif
+ && lookup_attribute ("omp declare target link", DECL_ATTRIBUTES (it));
- 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)));
+ 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
+ {
+#ifdef ACCEL_COMPILER
+ /* For "omp declare target link" vars add address of the pointer to
+ the target table, instead of address of the var. */
+ tree value_expr = DECL_VALUE_EXPR (it);
+ tree link_ptr_decl = TREE_OPERAND (value_expr, 0);
+ varpool_node::finalize_decl (link_ptr_decl);
+ addr = build_fold_addr_expr (link_ptr_decl);
+#else
+ addr = build_fold_addr_expr (it);
+#endif
+
+ /* Most significant bit of the size marks "omp declare target link"
+ vars in host and target tables. */
+ unsigned HOST_WIDE_INT isize = tree_to_uhwi (size);
+ isize |= 1ULL << (int_size_in_bytes (const_ptr_type_node)
+ * BITS_PER_UNIT - 1);
+ size = wide_int_to_tree (const_ptr_type_node, isize);
+ }
+
+ CONSTRUCTOR_APPEND_ELT (v_ctor, NULL_TREE, addr);
+ if (is_var)
+ CONSTRUCTOR_APPEND_ELT (v_ctor, NULL_TREE, size);
}
}
@@ -19831,4 +19866,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 43ce3d5..c72b38b 100644
--- a/gcc/passes.def
+++ b/gcc/passes.def
@@ -170,6 +170,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 e1cbce9..a13a865 100644
--- a/gcc/tree-pass.h
+++ b/gcc/tree-pass.h
@@ -417,6 +417,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 5e4fcbf..d0101a1 100644
--- a/gcc/varpool.c
+++ b/gcc/varpool.c
@@ -748,6 +748,13 @@ symbol_table::output_variables (void)
/* Handled in output_in_order. */
if (node->no_reorder)
continue;
+#ifdef ACCEL_COMPILER
+ /* Do not assemble "omp declare target link" vars. */
+ if (DECL_HAS_VALUE_EXPR_P (node->decl)
+ && lookup_attribute ("omp declare target link",
+ DECL_ATTRIBUTES (node->decl)))
+ continue;
+#endif
if (node->assemble_decl ())
changed = true;
}
diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h
index 9d9949f..73aa513 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 932b176..1ab30f7 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -464,7 +464,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
@@ -628,11 +628,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;
@@ -752,6 +760,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++;
}
}
@@ -884,6 +902,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
@@ -1020,31 +1041,40 @@ gomp_load_image_to_device (struct gomp_device_descr *devicep, unsigned version,
k->tgt_offset = target_table[i].start;
k->refcount = REFCOUNT_INFINITY;
k->async_refcount = 0;
+ k->link_key = NULL;
array->left = NULL;
array->right = NULL;
splay_tree_insert (&devicep->mem_map, array);
array++;
}
+ /* Most significant bit of the size in host and target tables marks
+ "omp declare target link" variables. */
+ const uintptr_t link_bit = 1ULL << (sizeof (uintptr_t) * __CHAR_BIT__ - 1);
+ const uintptr_t size_mask = ~link_bit;
+
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;
+
+ if ((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;
k->host_start = (uintptr_t) host_var_table[i * 2];
- k->host_end = k->host_start + (uintptr_t) host_var_table[i * 2 + 1];
+ k->host_end
+ = k->host_start + (size_mask & (uintptr_t) host_var_table[i * 2 + 1]);
k->tgt = tgt;
k->tgt_offset = target_var->start;
- k->refcount = REFCOUNT_INFINITY;
+ k->refcount = target_size & link_bit ? REFCOUNT_LINK : REFCOUNT_INFINITY;
k->async_refcount = 0;
+ k->link_key = NULL;
array->left = NULL;
array->right = NULL;
splay_tree_insert (&devicep->mem_map, array);
@@ -1072,7 +1102,6 @@ gomp_unload_image_from_device (struct gomp_device_descr *devicep,
int num_funcs = host_funcs_end - host_func_table;
int num_vars = (host_vars_end - host_var_table) / 2;
- unsigned j;
struct splay_tree_key_s k;
splay_tree_key node = NULL;
@@ -1088,21 +1117,46 @@ gomp_unload_image_from_device (struct gomp_device_descr *devicep,
devicep->unload_image_func (devicep->target_id, version, target_data);
/* Remove mappings from splay tree. */
- for (j = 0; j < num_funcs; j++)
+ int i;
+ for (i = 0; i < num_funcs; i++)
{
- k.host_start = (uintptr_t) host_func_table[j];
+ k.host_start = (uintptr_t) host_func_table[i];
k.host_end = k.host_start + 1;
splay_tree_remove (&devicep->mem_map, &k);
}
- for (j = 0; j < num_vars; j++)
+ /* Most significant bit of the size in host and target tables marks
+ "omp declare target link" variables. */
+ const uintptr_t link_bit = 1ULL << (sizeof (uintptr_t) * __CHAR_BIT__ - 1);
+ const uintptr_t size_mask = ~link_bit;
+ bool is_tgt_unmapped = false;
+
+ for (i = 0; i < num_vars; i++)
{
- k.host_start = (uintptr_t) host_var_table[j * 2];
- k.host_end = k.host_start + (uintptr_t) host_var_table[j * 2 + 1];
- splay_tree_remove (&devicep->mem_map, &k);
+ k.host_start = (uintptr_t) host_var_table[i * 2];
+ k.host_end
+ = k.host_start + (size_mask & (uintptr_t) host_var_table[i * 2 + 1]);
+
+ if (!(link_bit & (uintptr_t) host_var_table[i * 2 + 1]))
+ splay_tree_remove (&devicep->mem_map, &k);
+ else
+ {
+ splay_tree_key n = splay_tree_lookup (&devicep->mem_map, &k);
+ splay_tree_remove (&devicep->mem_map, n);
+ if (n->link_key)
+ {
+ if (n->tgt->refcount > 1)
+ n->tgt->refcount--;
+ else
+ {
+ is_tgt_unmapped = true;
+ gomp_unmap_tgt (n->tgt);
+ }
+ }
+ }
}
- if (node)
+ if (node && !is_tgt_unmapped)
{
free (node->tgt);
free (node);
@@ -1658,6 +1712,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
More information about the Gcc-patches
mailing list