diff mbox

[gomp4.5] Handle #pragma omp declare target link

Message ID 20151116154043.GA18854@msticlxl57.ims.intel.com
State New
Headers show

Commit Message

Ilya Verbin Nov. 16, 2015, 3:40 p.m. UTC 
Hi!

On Mon, Oct 26, 2015 at 20:49:40 +0100, Jakub Jelinek wrote:
> On Mon, Oct 26, 2015 at 10:39:04PM +0300, Ilya Verbin wrote:
> > > Without declare target link or to, you can't use the global variables
> > > in orphaned accelerated routines (unless you e.g. take the address of the
> > > mapped variable in the region and pass it around).
> > > The to variables (non-deferred) are always mapped and are initialized with
> > > the original initializer, refcount is infinity.  link (deferred) work more
> > > like the normal mapping, referencing those vars when they aren't explicitly
> > > (or implicitly) mapped is unspecified behavior, if it is e.g. mapped freshly
> > > with to kind, it gets the current value of the host var rather than the
> > > original one.  But, beyond the mapping the compiler needs to ensure that
> > > all uses of the link global var (or perhaps just all uses of the link global
> > > var outside of the target construct body where it is mapped, because you
> > > could use there the pointer you got from GOMP_target) are replaced by
> > > dereference of some artificial pointer, so a becomes *a_tmp and &a becomes
> > > &*a_tmp, and that the runtime library during registration of the tables is
> > > told about the address of this artificial pointer.  During registration,
> > > I'd expect it would stick an entry for this range into the table, with some
> > > special flag or something similar, indicating that it is deferred mapping
> > > and where the offloading device pointer is.  During mapping, it would map it
> > > as any other not yet mapped object, but additionally would also set this
> > > device pointer to the device address of the mapped object.  We also need to
> > > ensure that when we drop the refcount of that mapping back to 0, we get it
> > > back to the state where it is described as a range with registered deferred
> > > mapping and where the device pointer is.
> > 
> > Ok, got it, I'll try implement this...
> 
> Thanks.
> 
> > > > > we actually replace the variables with pointers to variables, then need
> > > > > to somehow also mark those in the offloading tables, so that the library
> > > > 
> > > > I see 2 possible options: use the MSB of the size, or introduce the third field
> > > > for flags.
> > > 
> > > Well, it can be either recorded in the host variable tables (which contain
> > > address and size pair, right), or in corresponding offloading device table
> > > (which contains the pointer, something else?).
> > 
> > It contains a size too, which is checked in libgomp:
> > 	  gomp_fatal ("Can't map target variables (size mismatch)");
> > Yes, we can remove this check, and use second field in device table for flags.
> 
> Yeah, or e.g. just use MSB of that size (so check that either the size is
> the same (then it is target to) or it is MSB | size (then it is target link).
> Objects larger than half of the address space aren't really supportable
> anyway.

Here is WIP patch, not for check-in.  There are still many FIXMEs, which I am
going to resolve, however target-link-1.c testcase pass.
Is this approach correct?  Any comments on FIXMEs?




  -- Ilya

Comments

Jakub Jelinek Nov. 19, 2015, 3:31 p.m. UTC | #1 
On Mon, Nov 16, 2015 at 06:40:43PM +0300, Ilya Verbin wrote:
> Here is WIP patch, not for check-in.  There are still many FIXMEs, which I am
> going to resolve, however target-link-1.c testcase pass.
> Is this approach correct?  Any comments on FIXMEs?
> 
> 
> diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
> index 23d0107..58771c0 100644
> --- a/gcc/c/c-parser.c
> +++ b/gcc/c/c-parser.c
> @@ -15895,7 +15895,10 @@ c_parser_omp_declare_target (c_parser *parser)
>  	      g->have_offload = true;
>  	      if (is_a <varpool_node *> (node))
>  		{
> -		  vec_safe_push (offload_vars, t);
> +		  omp_offload_var var;
> +		  var.decl = t;
> +		  var.link_ptr_decl = NULL_TREE;
> +		  vec_safe_push (offload_vars, var);
>  		  node->force_output = 1;
>  		}

Another possible approach would be to keep offload_vars as
vector of trees, and simply push 2 trees in each case.
Or not to change this at all, see below.

> @@ -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.

> +      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.

> +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.

> +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.

> +		  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.

> 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.

	Jakub
diff mbox

Patch

diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index 23d0107..58771c0 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -15895,7 +15895,10 @@  c_parser_omp_declare_target (c_parser *parser)
 	      g->have_offload = true;
 	      if (is_a <varpool_node *> (node))
 		{
-		  vec_safe_push (offload_vars, t);
+		  omp_offload_var var;
+		  var.decl = t;
+		  var.link_ptr_decl = NULL_TREE;
+		  vec_safe_push (offload_vars, var);
 		  node->force_output = 1;
 		}
 #endif
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index d1f4970..b890f6d 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -34999,7 +34999,10 @@  cp_parser_omp_declare_target (cp_parser *parser, cp_token *pragma_tok)
 	      g->have_offload = true;
 	      if (is_a <varpool_node *> (node))
 		{
-		  vec_safe_push (offload_vars, t);
+		  omp_offload_var var;
+		  var.decl = t;
+		  var.link_ptr_decl = NULL_TREE;
+		  vec_safe_push (offload_vars, var);
 		  node->force_output = 1;
 		}
 #endif
diff --git a/gcc/lto-cgraph.c b/gcc/lto-cgraph.c
index 67a9024..878a9c5 100644
--- a/gcc/lto-cgraph.c
+++ b/gcc/lto-cgraph.c
@@ -1106,7 +1106,7 @@  output_offload_tables (void)
       streamer_write_enum (ob->main_stream, LTO_symtab_tags,
 			   LTO_symtab_last_tag, LTO_symtab_variable);
       lto_output_var_decl_index (ob->decl_state, ob->main_stream,
-				 (*offload_vars)[i]);
+				 (*offload_vars)[i].decl);
     }
 
   streamer_write_uhwi_stream (ob->main_stream, 0);
@@ -1902,7 +1902,10 @@  input_offload_tables (void)
 	      int decl_index = streamer_read_uhwi (ib);
 	      tree var_decl
 		= lto_file_decl_data_get_var_decl (file_data, decl_index);
-	      vec_safe_push (offload_vars, var_decl);
+	      omp_offload_var var;
+	      var.decl = var_decl;
+	      var.link_ptr_decl = NULL_TREE;
+	      vec_safe_push (offload_vars, var);
 	    }
 	  else
 	    fatal_error (input_location,
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index ee33551..5900f1a 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -373,7 +373,8 @@  unshare_and_remap (tree x, tree from, tree to)
 }
 
 /* Holds offload tables with decls.  */
-vec<tree, va_gc> *offload_funcs, *offload_vars;
+vec<tree, va_gc> *offload_funcs;
+vec<omp_offload_var, va_gc> *offload_vars;
 
 /* Convenience function for calling scan_omp_1_op on tree operands.  */
 
@@ -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)))
 	    break;
 	  if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
 	      && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER)
@@ -18331,23 +18335,50 @@  make_pass_omp_simd_clone (gcc::context *ctxt)
   return new pass_omp_simd_clone (ctxt);
 }
 
-/* Helper function for omp_finish_file routine.  Takes decls from V_DECLS and
-   adds their addresses and sizes to constructor-vector V_CTOR.  */
+/* Helper function for omp_finish_file routine.  Takes func decls from V_DECLS
+   and adds their addresses to constructor-vector V_CTOR.  */
 static void
-add_decls_addresses_to_decl_constructor (vec<tree, va_gc> *v_decls,
-					 vec<constructor_elt, va_gc> *v_ctor)
+add_funcs_to_decl_constructor (vec<tree, va_gc> *v_decls,
+			       vec<constructor_elt, va_gc> *v_ctor)
 {
   unsigned len = vec_safe_length (v_decls);
   for (unsigned i = 0; i < len; i++)
     {
       tree it = (*v_decls)[i];
-      bool is_function = TREE_CODE (it) != VAR_DECL;
-
       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)));
+    }
+}
+
+/* Helper function for omp_finish_file routine.  Takes var decls from V_DECLS
+   and adds their addresses and sizes to constructor-vector V_CTOR.  */
+static void
+add_vars_to_decl_constructor (vec<omp_offload_var, va_gc> *v_decls,
+			      vec<constructor_elt, va_gc> *v_ctor)
+{
+  unsigned len = vec_safe_length (v_decls);
+  for (unsigned i = 0; i < len; i++)
+    {
+      omp_offload_var var = (*v_decls)[i];
+      tree addr;
+      tree size = fold_convert (const_ptr_type_node, DECL_SIZE_UNIT (var.decl));
+
+      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?  */
+	}
+
+      CONSTRUCTOR_APPEND_ELT (v_ctor, NULL_TREE, addr);
+      CONSTRUCTOR_APPEND_ELT (v_ctor, NULL_TREE, size);
     }
 }
 
@@ -18369,8 +18400,8 @@  omp_finish_file (void)
       vec_alloc (v_f, num_funcs);
       vec_alloc (v_v, num_vars * 2);
 
-      add_decls_addresses_to_decl_constructor (offload_funcs, v_f);
-      add_decls_addresses_to_decl_constructor (offload_vars, v_v);
+      add_funcs_to_decl_constructor (offload_funcs, v_f);
+      add_vars_to_decl_constructor (offload_vars, v_v);
 
       tree vars_decl_type = build_array_type_nelts (pointer_sized_int_node,
 						    num_vars * 2);
@@ -18412,7 +18443,7 @@  omp_finish_file (void)
 	}
       for (unsigned i = 0; i < num_vars; i++)
 	{
-	  tree it = (*offload_vars)[i];
+	  tree it = (*offload_vars)[i].decl;
 	  targetm.record_offload_symbol (it);
 	}
     }
@@ -19538,4 +19569,145 @@  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
+      /* FIXME: Replace globals in target regions too or not?  */
+      return lookup_attribute ("omp declare target",
+			       DECL_ATTRIBUTES (fun->decl));
+#else
+      (void) fun;
+      return false;
+#endif
+    }
+
+  virtual unsigned execute (function *);
+};
+
+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?  */
+
+	      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;
+		      }
+		  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);
+		      TREE_TYPE (ptr_decl) = ptype;
+		      DECL_MODE (ptr_decl) = TYPE_MODE (ptype);
+		      DECL_SIZE (ptr_decl) = TYPE_SIZE (ptype);
+		      DECL_SIZE_UNIT (ptr_decl) = TYPE_SIZE_UNIT (ptype);
+		      DECL_ARTIFICIAL (ptr_decl) = 1;
+		      /* FIXME: Add new function clone_variable_name?
+			 clone_function_name adds dots into the name, which are
+			 bad for vars.  */
+		      DECL_NAME (ptr_decl)
+			= clone_function_name (var, "linkptr");
+		      SET_DECL_ASSEMBLER_NAME (ptr_decl, DECL_NAME (ptr_decl));
+		      SET_DECL_RTL (ptr_decl, NULL);
+		      varpool_node::finalize_decl (ptr_decl);
+		      table_entry->link_ptr_decl = ptr_decl;
+		    }
+
+		  /* Replace the use of var with dereference of ptr_decl.  */
+		  tree tmp_ssa = make_temp_ssa_name (ptype, NULL, "linkptr");
+		  gimple *new_stmt = gimple_build_assign (tmp_ssa, ptr_decl);
+		  gsi_insert_before (&gsi, new_stmt, GSI_SAME_STMT);
+		  tree mem_ref = build_simple_mem_ref (tmp_ssa);
+
+		  if (TREE_CODE (op) == VAR_DECL)
+		    *gimple_op_ptr (stmt, i) = mem_ref;
+		  else if (TREE_CODE (op) == ADDR_EXPR)
+		    {
+		      tree op1 = TREE_OPERAND (op, 0);
+		      if (TREE_CODE (op1) == VAR_DECL)
+			TREE_OPERAND (op, 0) = mem_ref;
+		      recompute_tree_invariant_for_addr_expr (op);
+		    }
+		  update_stmt (stmt);
+		}
+	    }
+	}
+    }
+
+  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/omp-low.h b/gcc/omp-low.h
index ee0f8ac..c6e4d5a 100644
--- a/gcc/omp-low.h
+++ b/gcc/omp-low.h
@@ -34,7 +34,16 @@  extern tree get_oacc_fn_attrib (tree);
 extern int get_oacc_ifn_dim_arg (const gimple *);
 extern int get_oacc_fn_dim_size (tree, int);
 
+struct omp_offload_var
+{
+  /* Declaration representing global variable.  */
+  tree decl;
+
+  /* Artificial pointer for "omp declare target link" variables.  */
+  tree link_ptr_decl;
+};
+
 extern GTY(()) vec<tree, va_gc> *offload_funcs;
-extern GTY(()) vec<tree, va_gc> *offload_vars;
+extern GTY(()) vec<omp_offload_var, va_gc> *offload_vars;
 
 #endif /* GCC_OMP_LOW_H */
diff --git a/gcc/passes.def b/gcc/passes.def
index c0ab6b9..b32a5e5 100644
--- a/gcc/passes.def
+++ b/gcc/passes.def
@@ -151,6 +151,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 49e22a9..554f3d2 100644
--- a/gcc/tree-pass.h
+++ b/gcc/tree-pass.h
@@ -413,6 +413,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 478f365..ca8457d 100644
--- a/gcc/varpool.c
+++ b/gcc/varpool.c
@@ -156,7 +156,12 @@  varpool_node::get_create (tree decl)
 #ifdef ENABLE_OFFLOADING
       g->have_offload = true;
       if (!in_lto_p)
-	vec_safe_push (offload_vars, decl);
+	{
+	  omp_offload_var var;
+	  var.decl = decl;
+	  var.link_ptr_decl = NULL_TREE;
+	  vec_safe_push (offload_vars, var);
+	}
       node->force_output = 1;
 #endif
     }
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;
+
 /* Similar to gomp_realloc, but release register_lock before gomp_fatal.  */
 
 static void *
@@ -763,6 +774,21 @@  gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
 	}
     }
 
+  /* Set pointers to "omp declare target link" variables.  */
+  for (i = 0; i < mapnum; i++)
+    /* FIXME: Remove this ugly loop.  */
+    for (int j = 0; j < link_num; j++)
+      if (links[j].host_start == (uintptr_t) hostaddrs[i])
+	{
+	  cur_node.tgt_offset = gomp_map_val (tgt, hostaddrs, i);
+	  /* Set link pointer on target to the device address of the mapped
+	     object.  */
+	  devicep->host2dev_func (devicep->target_id,
+				  (void *) links[j].tgt_link_ptr,
+				  (void *) &cur_node.tgt_offset,
+				  sizeof (void *));
+	}
+
   /* If the variable from "omp target enter data" map-list was already mapped,
      tgt is not needed.  Otherwise tgt will be freed by gomp_unmap_vars or
      gomp_exit_data.  */
@@ -981,6 +1007,7 @@  gomp_load_image_to_device (struct gomp_device_descr *devicep, unsigned version,
 
   /* Insert host-target address mapping into splay tree.  */
   struct target_mem_desc *tgt = gomp_malloc (sizeof (*tgt));
+  /* FIXME: Do not allocate space for link vars.  */
   tgt->array = gomp_malloc ((num_funcs + num_vars) * sizeof (*tgt->array));
   tgt->refcount = REFCOUNT_INFINITY;
   tgt->tgt_start = 0;
@@ -1009,26 +1036,44 @@  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)
 	{
-	  gomp_mutex_unlock (&devicep->lock);
-	  if (is_register_lock)
-	    gomp_mutex_unlock (&register_lock);
-	  gomp_fatal ("Can't map target variables (size mismatch)");
-	}
+	  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)");
+	    }
 
-      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->tgt = tgt;
-      k->tgt_offset = target_var->start;
-      k->refcount = REFCOUNT_INFINITY;
-      k->async_refcount = 0;
-      array->left = NULL;
-      array->right = NULL;
-      splay_tree_insert (&devicep->mem_map, array);
-      array++;
+	  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->tgt = tgt;
+	  k->tgt_offset = target_var->start;
+	  k->refcount = REFCOUNT_INFINITY;
+	  k->async_refcount = 0;
+	  array->left = NULL;
+	  array->right = NULL;
+	  splay_tree_insert (&devicep->mem_map, array);
+	  array++;
+	}
+      else
+	{
+	  /* Do not map "omp declare target link" variables, only keep target
+	     address of the artificial pointer.  */
+	  /* FIXME: Where to keep it?  */
+	  struct link_struct l;
+	  l.host_start = (uintptr_t) host_var_table[i * 2];
+	  l.tgt_link_ptr = target_var->start;
+	  links[link_num++] = l;
+	}
     }
 
   free (target_table);
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..332bc14
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/target-link-1.c
@@ -0,0 +1,56 @@ 
+int a = 1, b = 1;
+double c = 1.0;
+long long d[27];
+#pragma omp declare target link (a) to (b) link (c, d)
+
+/* FIXME: When the function is inlined, it gets the wrong values.  */
+__attribute__((noinline, noclone)) int
+foo (void)
+{
+  return a++ + b++;
+}
+
+/* FIXME: When the function is inlined, it gets the wrong values.  */
+__attribute__((noinline, noclone)) int
+bar (void)
+{
+  int *p1 = &a;
+  int *p2 = &b;
+  c += 0.1;
+  d[10]++; /* FIXME: Support arrays in pass_omp_target_link::execute.  */
+  return *p1 + *p2;
+}
+
+#pragma omp declare target (foo, bar)
+
+int
+main ()
+{
+  int res;
+  a = b = 2;
+  #pragma omp target map (to: a, b, c, d) map (from: res)
+  {
+    a; c; d; /* FIXME: Do not remove map(a,c,d) during gimplification.  */
+    res = foo () + foo ();
+    res += bar ();
+  }
+
+  int shared_mem = 0;
+  #pragma omp target map (alloc: shared_mem)
+    shared_mem = 1;
+
+  if ((shared_mem && res != (2 + 2) + (3 + 3) + (4 + 4))
+      || (!shared_mem && res != (2 + 1) + (3 + 2) + (4 + 3)))
+    __builtin_abort ();
+
+  #pragma omp target map (to: a) map (from: res)
+  {
+    a; /* FIXME: Do not remove map(a) during gimplification.  */
+    res = foo ();
+  }
+
+  if ((shared_mem && res != 4 + 4) || (!shared_mem && res != 2 + 3))
+    __builtin_abort ();
+
+  return 0;
+}