[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

diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index a9dcbd8..d095a19 100644 (file)
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,3 +1,9 @@
+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.

diff --git a/libgomp/oacc-parallel.c b/libgomp/oacc-parallel.c
index b80ace5..1e08af7 100644 (file)
--- a/libgomp/oacc-parallel.c
+++ b/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);

diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr88288.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr88288.c
new file mode 100644 (file)
index 0000000..d13e335
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr88288.c
@@ -0,0 +1,41 @@
+/* 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;
+}