Re: [gomp4.1] Support #pragma omp target {enter,exit} data

[Ilya Verbin <iverbin at gmail dot com>]

On Thu, Jul 30, 2015 at 10:12:59 +0200, Jakub Jelinek wrote:
> On Wed, Jul 29, 2015 at 10:06:52PM +0300, Ilya Verbin wrote:
> > @@ -1241,6 +1245,62 @@ GOMP_target_update (int device, const void *unused, size_t mapnum,
> >    gomp_update (devicep, mapnum, hostaddrs, sizes, kinds, false);
> >  }
> >  
> > +static void
> > +gomp_exit_data (struct gomp_device_descr *devicep, size_t mapnum,
> > +		void **hostaddrs, size_t *sizes, unsigned short *kinds)
> > +{
> > +  const int typemask = 0xff;
> > +  size_t i;
> > +  gomp_mutex_lock (&devicep->lock);
> > +  for (i = 0; i < mapnum; i++)
> > +    {
> > +      struct splay_tree_key_s cur_node;
> > +      unsigned char kind = kinds[i] & typemask;
> > +      switch (kind)
> > +	{
> > +	case GOMP_MAP_FROM:
> > +	case GOMP_MAP_ALWAYS_FROM:
> > +	case GOMP_MAP_DELETE:
> > +	case GOMP_MAP_RELEASE:
> 
> Please handle here GOMP_MAP_ZERO_LEN_ARRAY_SECTION too.
> It should use gomp_map_lookup (while all others splay_tree_lookup),
> otherwise it is the same as GOMP_MAP_RELEASE.

Done.

> > @@ -1280,13 +1337,20 @@ GOMP_target_enter_exit_data (int device, size_t mapnum, void **hostaddrs,
> >      }
> >  
> >    if (is_enter_data)
> > -    {
> > -      /* TODO  */
> > -    }
> > +    for (i = 0; i < mapnum; i++)
> > +      {
> > +	struct target_mem_desc *tgt_var
> > +	  = gomp_map_vars (devicep, 1, &hostaddrs[i], NULL, &sizes[i],
> > +			   &kinds[i], true, false);
> > +	tgt_var->refcount--;
> > +
> > +	/* If the variable was already mapped, tgt_var is not needed.  Otherwise
> > +	   tgt_var will be freed by gomp_unmap_vars or gomp_exit_data.  */
> > +	if (tgt_var->refcount == 0)
> > +	  free (tgt_var);
> 
> This is racy, you don't hold the device lock here anymore, so you shouldn't
> decrease refcounts or test it etc.
> I think better would be to change the bool is_target argument to
> gomp_map_vars into an enum, and use 3 values there for now
> - GOMP_VARS_MAP_TARGET, GOMP_VARS_MAP_DATA, GOMP_VARS_MAP_ENTER_DATA or so,
> and for GOMP_VARS_MAP_ENTER_DATA perform the decrement of refcount and
> freeing if it is zero (but then also better return NULL).

Fixed.

> > diff --git a/libgomp/testsuite/libgomp.c/target-20.c b/libgomp/testsuite/libgomp.c/target-20.c
> > new file mode 100644
> > index 0000000..ec7e245
> > --- /dev/null
> > +++ b/libgomp/testsuite/libgomp.c/target-20.c
> > @@ -0,0 +1,111 @@
> > +/* { dg-require-effective-target offload_device } */
> 
> This test will fail on HSA, you don't assume just that it doesn't
> fallback to host, but also non-shared address space.

Fixed.

make check-target-libgomp passed.  ok?


libgomp/
	* libgomp.h (enum gomp_map_vars_kind): New.
	(gomp_map_vars): Change type of the argument from bool to enum
	gomp_map_vars_kind.
	* oacc-mem.c (acc_map_data, present_create_copy,
	gomp_acc_insert_pointer): Pass GOMP_MAP_VARS_OPENACC instead of false to
	gomp_map_vars.
	* oacc-parallel.c (GOACC_parallel, GOACC_data_start): Likewise.
	* target.c (gomp_map_vars_existing): Fix target address for 'always to'
	array sections.
	(gomp_map_vars): Change type of the argument from bool to enum
	gomp_map_vars_kind, fixup its usage.  Set tgt->refcount to 0 if called
	from GOMP_target_enter_exit_data.  Free tgt if called from
	GOMP_target_enter_exit_data and nothing has been mapped.
	(gomp_unmap_vars): Decrement k->refcount when it is 1 and
	k->async_refcount is 0.
	(gomp_offload_image_to_device): Set tgt's refcount to infinity.
	(GOMP_target, GOMP_target_41): Pass GOMP_MAP_VARS_TARGET instead of true
	to gomp_map_vars.
	(gomp_target_data_fallback, GOMP_target_data, GOMP_target_data_41): Pass
	GOMP_MAP_VARS_DATA instead of false to gomp_map_vars.
	(gomp_exit_data): New static function.
	(GOMP_target_enter_exit_data): Support mapping/unmapping.
	* testsuite/libgomp.c/target-11.c: Extend for testing 'always to' array
	sections.
	* testsuite/libgomp.c/target-20.c: New test.


diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h
index 707acaf..9031649 100644
--- a/libgomp/libgomp.h
+++ b/libgomp/libgomp.h
@@ -787,12 +787,22 @@ struct gomp_device_descr
   acc_dispatch_t openacc;
 };
 
+/* Kind of the pragma, for which gomp_map_vars () is called.  */
+enum gomp_map_vars_kind
+{
+  GOMP_MAP_VARS_OPENACC,
+  GOMP_MAP_VARS_TARGET,
+  GOMP_MAP_VARS_DATA,
+  GOMP_MAP_VARS_ENTER_DATA
+};
+
 extern void gomp_acc_insert_pointer (size_t, void **, size_t *, void *);
 extern void gomp_acc_remove_pointer (void *, bool, int, int);
 
 extern struct target_mem_desc *gomp_map_vars (struct gomp_device_descr *,
 					      size_t, void **, void **,
-					      size_t *, void *, bool, bool);
+					      size_t *, void *, bool,
+					      enum gomp_map_vars_kind);
 extern void gomp_copy_from_async (struct target_mem_desc *);
 extern void gomp_unmap_vars (struct target_mem_desc *, bool);
 extern void gomp_init_device (struct gomp_device_descr *);
diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c
index c0fcb07..af067d6 100644
--- a/libgomp/oacc-mem.c
+++ b/libgomp/oacc-mem.c
@@ -289,7 +289,8 @@ acc_map_data (void *h, void *d, size_t s)
       if (d != h)
         gomp_fatal ("cannot map data on shared-memory system");
 
-      tgt = gomp_map_vars (NULL, 0, NULL, NULL, NULL, NULL, true, false);
+      tgt = gomp_map_vars (NULL, 0, NULL, NULL, NULL, NULL, true,
+			   GOMP_MAP_VARS_OPENACC);
     }
   else
     {
@@ -318,7 +319,7 @@ acc_map_data (void *h, void *d, size_t s)
       gomp_mutex_unlock (&acc_dev->lock);
 
       tgt = gomp_map_vars (acc_dev, mapnum, &hostaddrs, &devaddrs, &sizes,
-			   &kinds, true, false);
+			   &kinds, true, GOMP_MAP_VARS_OPENACC);
     }
 
   gomp_mutex_lock (&acc_dev->lock);
@@ -447,7 +448,7 @@ present_create_copy (unsigned f, void *h, size_t s)
       gomp_mutex_unlock (&acc_dev->lock);
 
       tgt = gomp_map_vars (acc_dev, mapnum, &hostaddrs, NULL, &s, &kinds, true,
-			   false);
+			   GOMP_MAP_VARS_OPENACC);
 
       gomp_mutex_lock (&acc_dev->lock);
 
@@ -594,7 +595,7 @@ gomp_acc_insert_pointer (size_t mapnum, void **hostaddrs, size_t *sizes,
 
   gomp_debug (0, "  %s: prepare mappings\n", __FUNCTION__);
   tgt = gomp_map_vars (acc_dev, mapnum, hostaddrs,
-		       NULL, sizes, kinds, true, false);
+		       NULL, sizes, kinds, true, GOMP_MAP_VARS_OPENACC);
   gomp_debug (0, "  %s: mappings prepared\n", __FUNCTION__);
 
   gomp_mutex_lock (&acc_dev->lock);
diff --git a/libgomp/oacc-parallel.c b/libgomp/oacc-parallel.c
index 8ea3dd1..38c4770 100644
--- a/libgomp/oacc-parallel.c
+++ b/libgomp/oacc-parallel.c
@@ -131,7 +131,7 @@ GOACC_parallel (int device, void (*fn) (void *),
     tgt_fn = (void (*)) fn;
 
   tgt = gomp_map_vars (acc_dev, mapnum, hostaddrs, NULL, sizes, kinds, true,
-		       false);
+		       GOMP_MAP_VARS_OPENACC);
 
   devaddrs = gomp_alloca (sizeof (void *) * mapnum);
   for (i = 0; i < mapnum; i++)
@@ -178,7 +178,8 @@ GOACC_data_start (int device, size_t mapnum,
   if ((acc_dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
       || host_fallback)
     {
-      tgt = gomp_map_vars (NULL, 0, NULL, NULL, NULL, NULL, true, false);
+      tgt = gomp_map_vars (NULL, 0, NULL, NULL, NULL, NULL, true,
+			   GOMP_MAP_VARS_OPENACC);
       tgt->prev = thr->mapped_data;
       thr->mapped_data = tgt;
 
@@ -187,7 +188,7 @@ GOACC_data_start (int device, size_t mapnum,
 
   gomp_debug (0, "  %s: prepare mappings\n", __FUNCTION__);
   tgt = gomp_map_vars (acc_dev, mapnum, hostaddrs, NULL, sizes, kinds, true,
-		       false);
+		       GOMP_MAP_VARS_OPENACC);
   gomp_debug (0, "  %s: mappings prepared\n", __FUNCTION__);
   tgt->prev = thr->mapped_data;
   thr->mapped_data = tgt;
diff --git a/libgomp/target.c b/libgomp/target.c
index d7f4693..565982b 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -188,7 +188,8 @@ gomp_map_vars_existing (struct gomp_device_descr *devicep, splay_tree_key oldn,
 
   if (GOMP_MAP_ALWAYS_TO_P (kind))
     devicep->host2dev_func (devicep->target_id,
-			    (void *) (oldn->tgt->tgt_start + oldn->tgt_offset),
+			    (void *) (oldn->tgt->tgt_start + oldn->tgt_offset
+				      + newn->host_start - oldn->host_start),
 			    (void *) newn->host_start,
 			    newn->host_end - newn->host_start);
   if (oldn->refcount != REFCOUNT_INFINITY)
@@ -247,7 +248,7 @@ gomp_map_pointer (struct target_mem_desc *tgt, uintptr_t host_ptr,
 attribute_hidden struct target_mem_desc *
 gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
 	       void **hostaddrs, void **devaddrs, size_t *sizes, void *kinds,
-	       bool short_mapkind, bool is_target)
+	       bool short_mapkind, enum gomp_map_vars_kind pragma_kind)
 {
   size_t i, tgt_align, tgt_size, not_found_cnt = 0;
   bool has_firstprivate = false;
@@ -258,7 +259,7 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
   struct target_mem_desc *tgt
     = gomp_malloc (sizeof (*tgt) + sizeof (tgt->list[0]) * mapnum);
   tgt->list_count = mapnum;
-  tgt->refcount = 1;
+  tgt->refcount = pragma_kind == GOMP_MAP_VARS_ENTER_DATA ? 0 : 1;
   tgt->device_descr = devicep;
 
   if (mapnum == 0)
@@ -266,7 +267,7 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
 
   tgt_align = sizeof (void *);
   tgt_size = 0;
-  if (is_target)
+  if (pragma_kind == GOMP_MAP_VARS_TARGET)
     {
       size_t align = 4 * sizeof (void *);
       tgt_align = align;
@@ -377,7 +378,7 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
       tgt->tgt_start = (uintptr_t) tgt->to_free;
       tgt->tgt_end = tgt->tgt_start + sizes[0];
     }
-  else if (not_found_cnt || is_target)
+  else if (not_found_cnt || pragma_kind == GOMP_MAP_VARS_TARGET)
     {
       /* Allocate tgt_align aligned tgt_size block of memory.  */
       /* FIXME: Perhaps change interface to allocate properly aligned
@@ -396,7 +397,7 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
     }
 
   tgt_size = 0;
-  if (is_target)
+  if (pragma_kind == GOMP_MAP_VARS_TARGET)
     tgt_size = mapnum * sizeof (void *);
 
   tgt->array = NULL;
@@ -560,7 +561,7 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
 	  }
     }
 
-  if (is_target)
+  if (pragma_kind == GOMP_MAP_VARS_TARGET)
     {
       for (i = 0; i < mapnum; i++)
 	{
@@ -587,6 +588,15 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
 	}
     }
 
+  /* 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.  */
+  if (pragma_kind == GOMP_MAP_VARS_ENTER_DATA && tgt->refcount == 0)
+    {
+      free (tgt);
+      tgt = NULL;
+    }
+
   gomp_mutex_unlock (&devicep->lock);
   return tgt;
 }
@@ -661,15 +671,18 @@ gomp_unmap_vars (struct target_mem_desc *tgt, bool do_copyfrom)
 	continue;
 
       bool do_unmap = false;
-      if (k->refcount > 1)
+      if (k->refcount > 1 && k->refcount != REFCOUNT_INFINITY)
+	k->refcount--;
+      else if (k->refcount == 1)
 	{
-	  if (k->refcount != REFCOUNT_INFINITY)
-	    k->refcount--;
+	  if (k->async_refcount > 0)
+	    k->async_refcount--;
+	  else
+	    {
+	      k->refcount--;
+	      do_unmap = true;
+	    }
 	}
-      else if (k->async_refcount > 0)
-	k->async_refcount--;
-      else
-	do_unmap = true;
 
       if ((do_unmap && do_copyfrom && tgt->list[i].copy_from)
 	  || tgt->list[i].always_copy_from)
@@ -794,7 +807,7 @@ gomp_offload_image_to_device (struct gomp_device_descr *devicep,
   /* Insert host-target address mapping into splay tree.  */
   struct target_mem_desc *tgt = gomp_malloc (sizeof (*tgt));
   tgt->array = gomp_malloc ((num_funcs + num_vars) * sizeof (*tgt->array));
-  tgt->refcount = 1;
+  tgt->refcount = REFCOUNT_INFINITY;
   tgt->tgt_start = 0;
   tgt->tgt_end = 0;
   tgt->to_free = NULL;
@@ -1080,7 +1093,7 @@ GOMP_target (int device, void (*fn) (void *), const void *unused,
 
   struct target_mem_desc *tgt_vars
     = gomp_map_vars (devicep, mapnum, hostaddrs, NULL, sizes, kinds, false,
-		     true);
+		     GOMP_MAP_VARS_TARGET);
   struct gomp_thread old_thr, *thr = gomp_thread ();
   old_thr = *thr;
   memset (thr, '\0', sizeof (*thr));
@@ -1140,7 +1153,7 @@ GOMP_target_41 (int device, void (*fn) (void *), size_t mapnum,
 
   struct target_mem_desc *tgt_vars
     = gomp_map_vars (devicep, mapnum, hostaddrs, NULL, sizes, kinds, true,
-		     true);
+		     GOMP_MAP_VARS_TARGET);
   struct gomp_thread old_thr, *thr = gomp_thread ();
   old_thr = *thr;
   memset (thr, '\0', sizeof (*thr));
@@ -1168,7 +1181,8 @@ gomp_target_data_fallback (void)
          new #pragma omp target data, otherwise GOMP_target_end_data
          would get out of sync.  */
       struct target_mem_desc *tgt
-	= gomp_map_vars (NULL, 0, NULL, NULL, NULL, NULL, false, false);
+	= gomp_map_vars (NULL, 0, NULL, NULL, NULL, NULL, false,
+			 GOMP_MAP_VARS_DATA);
       tgt->prev = icv->target_data;
       icv->target_data = tgt;
     }
@@ -1186,7 +1200,7 @@ GOMP_target_data (int device, const void *unused, size_t mapnum,
 
   struct target_mem_desc *tgt
     = gomp_map_vars (devicep, mapnum, hostaddrs, NULL, sizes, kinds, false,
-		     false);
+		     GOMP_MAP_VARS_DATA);
   struct gomp_task_icv *icv = gomp_icv (true);
   tgt->prev = icv->target_data;
   icv->target_data = tgt;
@@ -1204,7 +1218,7 @@ GOMP_target_data_41 (int device, size_t mapnum, void **hostaddrs, size_t *sizes,
 
   struct target_mem_desc *tgt
     = gomp_map_vars (devicep, mapnum, hostaddrs, NULL, sizes, kinds, true,
-		     false);
+		     GOMP_MAP_VARS_DATA);
   struct gomp_task_icv *icv = gomp_icv (true);
   tgt->prev = icv->target_data;
   icv->target_data = tgt;
@@ -1235,6 +1249,65 @@ GOMP_target_update (int device, const void *unused, size_t mapnum,
   gomp_update (devicep, mapnum, hostaddrs, sizes, kinds, false);
 }
 
+static void
+gomp_exit_data (struct gomp_device_descr *devicep, size_t mapnum,
+		void **hostaddrs, size_t *sizes, unsigned short *kinds)
+{
+  const int typemask = 0xff;
+  size_t i;
+  gomp_mutex_lock (&devicep->lock);
+  for (i = 0; i < mapnum; i++)
+    {
+      struct splay_tree_key_s cur_node;
+      unsigned char kind = kinds[i] & typemask;
+      switch (kind)
+	{
+	case GOMP_MAP_FROM:
+	case GOMP_MAP_ALWAYS_FROM:
+	case GOMP_MAP_DELETE:
+	case GOMP_MAP_RELEASE:
+	case GOMP_MAP_ZERO_LEN_ARRAY_SECTION:
+	  cur_node.host_start = (uintptr_t) hostaddrs[i];
+	  cur_node.host_end = cur_node.host_start + sizes[i];
+	  splay_tree_key k = kind == GOMP_MAP_ZERO_LEN_ARRAY_SECTION
+	    ? gomp_map_lookup (&devicep->mem_map, &cur_node)
+	    : splay_tree_lookup (&devicep->mem_map, &cur_node);
+	  if (!k)
+	    continue;
+
+	  if (k->refcount > 0 && k->refcount != REFCOUNT_INFINITY)
+	    k->refcount--;
+	  if (kind == GOMP_MAP_DELETE && k->refcount != REFCOUNT_INFINITY)
+	    k->refcount = 0;
+
+	  if ((kind == GOMP_MAP_FROM && k->refcount == 0)
+	      || kind == GOMP_MAP_ALWAYS_FROM)
+	    devicep->dev2host_func (devicep->target_id,
+				    (void *) cur_node.host_start,
+				    (void *) (k->tgt->tgt_start + k->tgt_offset
+					      + cur_node.host_start
+					      - k->host_start),
+				    cur_node.host_end - cur_node.host_start);
+	  if (k->refcount == 0)
+	    {
+	      splay_tree_remove (&devicep->mem_map, k);
+	      if (k->tgt->refcount > 1)
+		k->tgt->refcount--;
+	      else
+		gomp_unmap_tgt (k->tgt);
+	    }
+
+	  break;
+	default:
+	  gomp_mutex_unlock (&devicep->lock);
+	  gomp_fatal ("GOMP_target_enter_exit_data unhandled kind 0x%.2x",
+		      kind);
+	}
+    }
+
+  gomp_mutex_unlock (&devicep->lock);
+}
+
 void
 GOMP_target_enter_exit_data (int device, size_t mapnum, void **hostaddrs,
 			     size_t *sizes, unsigned short *kinds)
@@ -1253,9 +1326,6 @@ GOMP_target_enter_exit_data (int device, size_t mapnum, void **hostaddrs,
     {
       unsigned char kind = kinds[i] & typemask;
 
-      if (kind == GOMP_MAP_POINTER || kind == GOMP_MAP_TO_PSET)
-	continue;
-
       if (kind == GOMP_MAP_ALLOC
 	  || kind == GOMP_MAP_TO
 	  || kind == GOMP_MAP_ALWAYS_TO)
@@ -1267,20 +1337,19 @@ GOMP_target_enter_exit_data (int device, size_t mapnum, void **hostaddrs,
       if (kind == GOMP_MAP_FROM
 	  || kind == GOMP_MAP_ALWAYS_FROM
 	  || kind == GOMP_MAP_DELETE
-	  || kind == GOMP_MAP_RELEASE)
+	  || kind == GOMP_MAP_RELEASE
+	  || kind == GOMP_MAP_ZERO_LEN_ARRAY_SECTION)
 	break;
 
       gomp_fatal ("GOMP_target_enter_exit_data unhandled kind 0x%.2x", kind);
     }
 
   if (is_enter_data)
-    {
-      /* TODO  */
-    }
+    for (i = 0; i < mapnum; i++)
+      gomp_map_vars (devicep, 1, &hostaddrs[i], NULL, &sizes[i], &kinds[i],
+		     true, GOMP_MAP_VARS_ENTER_DATA);
   else
-    {
-      /* TODO  */
-    }
+    gomp_exit_data (devicep, mapnum, hostaddrs, sizes, kinds);
 }
 
 void
diff --git a/libgomp/testsuite/libgomp.c/target-11.c b/libgomp/testsuite/libgomp.c/target-11.c
index ed6a17a..625c286 100644
--- a/libgomp/testsuite/libgomp.c/target-11.c
+++ b/libgomp/testsuite/libgomp.c/target-11.c
@@ -9,6 +9,17 @@ void test_array_section (int *p)
 {
   #pragma omp target data map(alloc: p[0:N])
     {
+      int ok = 1;
+      for (int i = 10; i < 10 + 4; i++)
+	p[i] = 997 * i;
+
+      #pragma omp target map(always to:p[10:4]) map(tofrom: ok)
+	for (int i = 10; i < 10 + 4; i++)
+	  if (p[i] != 997 * i)
+	    ok = 0;
+
+      assert (ok);
+
       #pragma omp target map(always from:p[7:9])
 	for (int i = 0; i < N; i++)
 	  p[i] = i;
diff --git a/libgomp/testsuite/libgomp.c/target-20.c b/libgomp/testsuite/libgomp.c/target-20.c
new file mode 100644
index 0000000..3f4e798
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/target-20.c
@@ -0,0 +1,120 @@
+/* { dg-require-effective-target offload_device_nonshared_as } */
+
+#include <stdlib.h>
+#include <assert.h>
+
+#define N 40
+
+int sum;
+int var1 = 1;
+int var2 = 2;
+
+#pragma omp declare target
+int D[N];
+#pragma omp end declare target
+
+void enter_data (int *X)
+{
+  #pragma omp target enter data map(to: var1, var2, X[:N]) map(alloc: sum)
+}
+
+void exit_data_0 (int *D)
+{
+  #pragma omp target exit data map(delete: D[:N])
+}
+
+void exit_data_1 ()
+{
+  #pragma omp target exit data map(from: var1)
+}
+
+void exit_data_2 (int *X)
+{
+  #pragma omp target exit data map(from: var2) map(release: X[:N], sum)
+}
+
+void exit_data_3 (int *p)
+{
+  #pragma omp target exit data map(from: p[:0])
+}
+
+void test_nested ()
+{
+  int X = 0, Y = 0, Z = 0;
+
+  #pragma omp target data map(from: X, Y, Z)
+    {
+      #pragma omp target data map(from: X, Y, Z)
+	{
+	  #pragma omp target map(from: X, Y, Z)
+	    X = Y = Z = 1337;
+	  assert (X == 0);
+	  assert (Y == 0);
+	  assert (Z == 0);
+
+	  #pragma omp target exit data map(from: X) map(release: Y)
+	  assert (X == 0);
+	  assert (Y == 0);
+
+	  #pragma omp target exit data map(release: Y) map(delete: Z)
+	  assert (Y == 0);
+	  assert (Z == 0);
+	}
+      assert (X == 1337);
+      assert (Y == 0);
+      assert (Z == 0);
+
+      #pragma omp target map(from: X)
+	X = 2448;
+      assert (X == 2448);
+      assert (Y == 0);
+      assert (Z == 0);
+
+      X = 4896;
+    }
+  assert (X == 4896);
+  assert (Y == 0);
+  assert (Z == 0);
+}
+
+int main ()
+{
+  int *X = malloc (N * sizeof (int));
+  int *Y = malloc (N * sizeof (int));
+  X[10] = 10;
+  Y[20] = 20;
+  enter_data (X);
+
+  exit_data_0 (D); /* This should have no effect on D.  */
+
+  #pragma omp target map(alloc: var1, var2, X[:N]) map(to: Y[:N]) map(always from: sum)
+    {
+      var1 += X[10];
+      var2 += Y[20];
+      sum = var1 + var2;
+      D[sum]++;
+    }
+
+  assert (var1 == 1);
+  assert (var2 == 2);
+  assert (sum == 33);
+
+  exit_data_1 ();
+  assert (var1 == 11);
+  assert (var2 == 2);
+
+  /* Increase refcount of already mapped X[0:N].  */
+  #pragma omp target enter data map(alloc: X[16:1])
+
+  exit_data_2 (X);
+  assert (var2 == 22);
+
+  exit_data_3 (X + 5); /* Unmap X[0:N].  */
+
+  free (X);
+  free (Y);
+
+  test_nested ();
+
+  return 0;
+}


  -- Ilya