Re: [gomp4.1] Add new versions of GOMP_target{,_data,_update} and GOMP_target_enter_exit_data

[Ilya Verbin <iverbin at gmail dot com>]

On Wed, Jun 24, 2015 at 13:39:03 +0200, Jakub Jelinek wrote:
> Thinking about this more, for always modifier this isn't really sufficient.
> Consider:
> void
> foo (int *p)
> {
>   #pragma omp target data (alloc:p[0:32])
>   {
>     #pragma omp target data (always, from:p[7:9])
>     {
>       ...
>     }
>   }
> }
> If all we record is the corresponding splay_tree and the flags
> (from/always_from), then this would try to copy from the device
> the whole array section, rather than just the small portion of it.
> So, supposedly in addition to the splay_tree for always from case we also
> need to remember e.g. [relative offset, length] within the splay tree
> object.

Indeed, here is the fix, make check-target-libgomp passed.


libgomp/
	* libgomp.h (struct target_var_desc): Add offset and length.
	* target.c (gomp_map_vars_existing): New argument tgt_var, fill it.
	(gomp_map_vars): Move filling of tgt->list[i] into
	gomp_map_vars_existing.  Add missed case GOMP_MAP_ALWAYS_FROM.
	(gomp_unmap_vars): Add list[i].offset to host and target addresses,
	use list[i].length instead of k->host_end - k->host_start.
	* testsuite/libgomp.c/target-11.c: Extend for testing array sections.


diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h
index bd17828..c48e708 100644
--- a/libgomp/libgomp.h
+++ b/libgomp/libgomp.h
@@ -644,6 +644,12 @@ struct target_var_desc {
   bool copy_from;
   /* True if data always should be copied from device to host at the end.  */
   bool always_copy_from;
+  /* Used for unmapping of array sections, can be nonzero only when
+     always_copy_from is true.  */
+  uintptr_t offset;
+  /* Used for unmapping of array sections, can be less than the size of the
+     whole object only when always_copy_from is true.  */
+  uintptr_t length;
 };
 
 struct target_mem_desc {
diff --git a/libgomp/target.c b/libgomp/target.c
index b1640c1..a394e95 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -149,8 +149,15 @@ resolve_device (int device_id)
 
 static inline void
 gomp_map_vars_existing (struct gomp_device_descr *devicep, splay_tree_key oldn,
-			splay_tree_key newn, unsigned char kind)
+			splay_tree_key newn, struct target_var_desc *tgt_var,
+			unsigned char kind)
 {
+  tgt_var->key = oldn;
+  tgt_var->copy_from = GOMP_MAP_COPY_FROM_P (kind);
+  tgt_var->always_copy_from = GOMP_MAP_ALWAYS_FROM_P (kind);
+  tgt_var->offset = newn->host_start - oldn->host_start;
+  tgt_var->length = newn->host_end - newn->host_start;
+
   if ((kind & GOMP_MAP_FLAG_FORCE)
       || oldn->host_start > newn->host_start
       || oldn->host_end < newn->host_end)
@@ -276,13 +283,8 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
 	cur_node.host_end = cur_node.host_start + sizeof (void *);
       splay_tree_key n = splay_tree_lookup (mem_map, &cur_node);
       if (n)
-	{
-	  tgt->list[i].key = n;
-	  tgt->list[i].copy_from = GOMP_MAP_COPY_FROM_P (kind & typemask);
-	  tgt->list[i].always_copy_from
-	    = GOMP_MAP_ALWAYS_FROM_P (kind & typemask);
-	  gomp_map_vars_existing (devicep, n, &cur_node, kind & typemask);
-	}
+	gomp_map_vars_existing (devicep, n, &cur_node, &tgt->list[i],
+				kind & typemask);
       else
 	{
 	  tgt->list[i].key = NULL;
@@ -367,13 +369,8 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
 	      k->host_end = k->host_start + sizeof (void *);
 	    splay_tree_key n = splay_tree_lookup (mem_map, k);
 	    if (n)
-	      {
-		tgt->list[i].key = n;
-		tgt->list[i].copy_from = GOMP_MAP_COPY_FROM_P (kind & typemask);
-		tgt->list[i].always_copy_from
-		  = GOMP_MAP_ALWAYS_FROM_P (kind & typemask);
-		gomp_map_vars_existing (devicep, n, k, kind & typemask);
-	      }
+	      gomp_map_vars_existing (devicep, n, k, &tgt->list[i],
+				      kind & typemask);
 	    else
 	      {
 		size_t align = (size_t) 1 << (kind >> rshift);
@@ -385,6 +382,8 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
 		tgt->list[i].copy_from = GOMP_MAP_COPY_FROM_P (kind & typemask);
 		tgt->list[i].always_copy_from
 		  = GOMP_MAP_ALWAYS_FROM_P (kind & typemask);
+		tgt->list[i].offset = 0;
+		tgt->list[i].length = k->host_end - k->host_start;
 		k->refcount = 1;
 		k->async_refcount = 0;
 		tgt->refcount++;
@@ -397,6 +396,7 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
 		  case GOMP_MAP_FROM:
 		  case GOMP_MAP_FORCE_ALLOC:
 		  case GOMP_MAP_FORCE_FROM:
+		  case GOMP_MAP_ALWAYS_FROM:
 		    break;
 		  case GOMP_MAP_TO:
 		  case GOMP_MAP_TOFROM:
@@ -587,9 +587,11 @@ gomp_unmap_vars (struct target_mem_desc *tgt, bool do_copyfrom)
 
       if ((do_unmap && do_copyfrom && tgt->list[i].copy_from)
 	  || tgt->list[i].always_copy_from)
-	devicep->dev2host_func (devicep->target_id, (void *) k->host_start,
-				(void *) (k->tgt->tgt_start + k->tgt_offset),
-				k->host_end - k->host_start);
+	devicep->dev2host_func (devicep->target_id,
+				(void *) (k->host_start + tgt->list[i].offset),
+				(void *) (k->tgt->tgt_start + k->tgt_offset
+					  + tgt->list[i].offset),
+				tgt->list[i].length);
       if (do_unmap)
 	{
 	  splay_tree_remove (&devicep->mem_map, k);
diff --git a/libgomp/testsuite/libgomp.c/target-11.c b/libgomp/testsuite/libgomp.c/target-11.c
index 0fd183b..b86097a 100644
--- a/libgomp/testsuite/libgomp.c/target-11.c
+++ b/libgomp/testsuite/libgomp.c/target-11.c
@@ -1,7 +1,20 @@
 /* { dg-require-effective-target offload_device } */
 
+#include <stdlib.h>
 #include <assert.h>
 
+#define N 32
+
+void test_array_section (int *p)
+{
+  #pragma omp target data map(alloc: p[0:N])
+    {
+      #pragma omp target map(always from:p[7:9])
+	for (int i = 0; i < N; i++)
+	  p[i] = i;
+    }
+}
+
 int main ()
 {
   int aa = 0, bb = 0, cc = 0, dd = 0;
@@ -47,5 +60,16 @@ int main ()
   assert (cc == 4);
   assert (dd == 4);
 
+  int *array = calloc (N, sizeof (int));
+  test_array_section (array);
+
+  for (int i = 0; i < 7; i++)
+    assert (array[i] == 0);
+  for (int i = 7; i < 7 + 9; i++)
+    assert (array[i] == i);
+  for (int i = 7 + 9; i < N; i++)
+    assert (array[i] == 0);
+
+  free (array);
   return 0;
 }


  -- Ilya