[PR88288, OpenACC, libgomp] Adjust offsets for present data clauses

Make libgomp respect the on device offset of subarrays which may arise in
present data clauses.

	libgomp/
	PR libgomp/88288
	* oacc-parallel.c (GOACC_parallel_keyed): Add offset to devaddrs.
	* testsuite/libgomp.oacc-c-c++-common/pr88288.c: New test.

Reviewed-by: Thomas Schwinge <thomas@codesourcery.com>

From-SVN: r266688

libgomp/ChangeLog

+2018-11-30  Cesar Philippidis  <cesar@codesourcery.com>
+
+	PR libgomp/88288
+	* oacc-parallel.c (GOACC_parallel_keyed): Add offset to devaddrs.
+	* testsuite/libgomp.oacc-c-c++-common/pr88288.c: New test.
+
 2018-11-30  Thomas Schwinge  <thomas@codesourcery.com>
 
 	* testsuite/libgomp.oacc-fortran/lib-16-2.f90: New file.

libgomp/oacc-parallel.c

@@ -232,7 +232,8 @@ GOACC_parallel_keyed (int device, void (*fn) (void *),
   devaddrs = gomp_alloca (sizeof (void *) * mapnum);
   for (i = 0; i < mapnum; i++)
     devaddrs[i] = (void *) (tgt->list[i].key->tgt->tgt_start
-			    + tgt->list[i].key->tgt_offset);
+			    + tgt->list[i].key->tgt_offset
+			    + tgt->list[i].offset);
 
   acc_dev->openacc.exec_func (tgt_fn, mapnum, hostaddrs, devaddrs,
 			      async, dims, tgt);

libgomp/testsuite/libgomp.oacc-c-c++-common/pr88288.c

+/* Test present data clauses in acc offloaded regions when the
+   subarray inside the present clause does not have the same base
+   offset value as the subarray in the enclosing acc data or acc enter
+   data variable.  */
+
+#include <assert.h>
+
+void
+offset (int *data, int n)
+{
+  int i;
+
+#pragma acc parallel loop present (data[0:n])
+  for (i = 0; i < n; i++)
+    data[i] = n;
+}
+
+int
+main ()
+{
+  const int n = 30;
+  int data[n], i;
+
+  for (i = 0; i < n; i++)
+    data[i] = -1;
+
+#pragma acc data copy(data[0:n])
+  {
+    offset (data + 10, 10);
+  }
+
+  for (i = 0; i < n; i++)
+    {
+      if (i < 10 || i >= 20)
+	assert (data[i] == -1);
+      else
+	assert (data[i] == 10);
+    }
+
+  return 0;
+}