From: Cesar Philippidis <cesar@codesourcery.com>
Date: Fri, 30 Nov 2018 20:39:49 +0000 (-0800)
Subject: [PR88288, OpenACC, libgomp] Adjust offsets for present data clauses
X-Git-Url: https://git.libre-soc.org/?a=commitdiff_plain;h=fe570ff8d4347ea98108da3cf0d4f3800294d5c5;p=gcc.git

[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 a9dcbd80820..d095a196fb6 100644
--- 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 b80ace58590..1e08af70b4d 100644
--- 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
index 00000000000..d13e3359a3e
--- /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;
+}