aboutsummaryrefslogtreecommitdiff
diff options
context:
space:
mode:
Diffstat
-rw-r--r--gcc/ChangeLog20
-rw-r--r--gcc/c-family/ChangeLog5
-rw-r--r--gcc/c-family/c-common.c2
-rw-r--r--gcc/cgraphunit.c7
-rw-r--r--gcc/gimplify.c4
-rw-r--r--gcc/lto/ChangeLog6
-rw-r--r--gcc/lto/lto.c35
-rw-r--r--gcc/omp-low.c131
-rw-r--r--gcc/passes.def1
-rw-r--r--gcc/tree-pass.h1
-rw-r--r--gcc/varpool.c7
-rw-r--r--libgomp/ChangeLog18
-rw-r--r--libgomp/libgomp.h5
-rw-r--r--libgomp/target.c87
-rw-r--r--libgomp/testsuite/libgomp.c/target-link-1.c63
15 files changed, 368 insertions, 24 deletions
diff --git a/gcc/ChangeLog b/gcc/ChangeLog
index 69fe3a7..02c2117 100644
--- a/gcc/ChangeLog
+++ b/gcc/ChangeLog
@@ -1,3 +1,23 @@
+2015-12-15 Ilya Verbin <ilya.verbin@intel.com>
+
+ * 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.
+ * 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.
+
2015-12-15 Bernd Schmidt <bschmidt@redhat.com>
PR middle-end/21273
diff --git a/gcc/c-family/ChangeLog b/gcc/c-family/ChangeLog
index db9a279..525cc16 100644
--- a/gcc/c-family/ChangeLog
+++ b/gcc/c-family/ChangeLog
@@ -1,3 +1,8 @@
+2015-12-15 Ilya Verbin <ilya.verbin@intel.com>
+
+ * c-common.c (c_common_attribute_table): Handle "omp declare target
+ link" attribute.
+
2015-12-14 Jakub Jelinek <jakub@redhat.com>
PR c/68833
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/ChangeLog b/gcc/lto/ChangeLog
index 6e90527..ac20a3f 100644
--- a/gcc/lto/ChangeLog
+++ b/gcc/lto/ChangeLog
@@ -1,3 +1,9 @@
+2015-12-15 Ilya Verbin <ilya.verbin@intel.com>
+
+ * lto.c: Include stringpool.h and fold-const.h.
+ (offload_handle_link_vars): New static function.
+ (lto_main): Call offload_handle_link_vars.
+
2015-12-10 Jan Hubicka <hubicka@ucw.cz>
* lto.c (lto_read_in_decl_state): Unpickle compressed bit.
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/ChangeLog b/libgomp/ChangeLog
index 8745927..9315d8b 100644
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,3 +1,21 @@
+2015-12-15 Ilya Verbin <ilya.verbin@intel.com>
+
+ * 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.
+
2015-12-14 Ilya Verbin <ilya.verbin@intel.com>
* libgomp.h (gomp_device_state): New enum.
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 (&register_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;
+}
generated by cgit v1.1 at 2025-08-09 08:24:38 +0000