This is another old gomp4 patch that corrects a bug where the runtime
was passing the wrong offset for subarray data to the accelerator. The
original description of this patch can be found here
<https://gcc.gnu.org/ml/gcc-patches/2016-08/msg01676.html>
I bootstrapped and regtested on x86_64/nvptx. Is it OK for trunk?
Thanks,
Cesar
>From fb743d8a45193c177cb0082400d140949e8c1e6d Mon Sep 17 00:00:00 2001
From: Cesar Philippidis <ce...@codesourcery.com>
Date: Wed, 24 Aug 2016 00:02:50 +0000
Subject: [PATCH 5/5] [libgomp, OpenACC] Adjust offsets for present data
clauses
2018-XX-YY Cesar Philippidis <ce...@codesourcery.com>
libgomp/
* oacc-parallel.c (GOACC_parallel_keyed): Add offset to devaddrs.
* testsuite/libgomp.oacc-c-c++-common/data_offset.c: New test.
* testsuite/libgomp.oacc-fortran/data_offset.f90: New test.
(cherry picked from gomp-4_0-branch r239723, 00c2585)
---
libgomp/oacc-parallel.c | 10 ++++-
.../libgomp.oacc-c-c++-common/data_offset.c | 41 ++++++++++++++++++
.../libgomp.oacc-fortran/data_offset.f90 | 43 +++++++++++++++++++
3 files changed, 92 insertions(+), 2 deletions(-)
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/data_offset.c
create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/data_offset.f90
diff --git a/libgomp/oacc-parallel.c b/libgomp/oacc-parallel.c
index b80ace58590..20e9ab2e251 100644
--- a/libgomp/oacc-parallel.c
+++ b/libgomp/oacc-parallel.c
@@ -231,8 +231,14 @@ 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);
+ {
+ if (tgt->list[i].key != NULL)
+ devaddrs[i] = (void *) (tgt->list[i].key->tgt->tgt_start
+ + tgt->list[i].key->tgt_offset
+ + tgt->list[i].offset);
+ else
+ devaddrs[i] = NULL;
+ }
acc_dev->openacc.exec_func (tgt_fn, mapnum, hostaddrs, devaddrs,
async, dims, tgt);
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/data_offset.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/data_offset.c
new file mode 100644
index 00000000000..ccbbfcab87b
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/data_offset.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;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/data_offset.f90 b/libgomp/testsuite/libgomp.oacc-fortran/data_offset.f90
new file mode 100644
index 00000000000..ff8ee39f964
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/data_offset.f90
@@ -0,0 +1,43 @@
+! 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.
+
+program test
+ implicit none
+
+ integer, parameter :: n = 30, m = 10
+ integer :: i
+ integer, allocatable :: data(:)
+ logical bounded
+
+ allocate (data(n))
+
+ data(:) = -1
+
+ !$acc data copy (data(5:20))
+ call test_data (data, n, m)
+ !$acc end data
+
+ do i = 1, n
+ bounded = i < m .or. i >= m+m
+ if (bounded .and. (data(i) /= -1)) then
+ call abort
+ else if (.not. bounded .and. data(i) /= 10) then
+ call abort
+ end if
+ end do
+
+ deallocate (data)
+end program test
+
+subroutine test_data (data, n, m)
+ implicit none
+
+ integer :: n, m, data(n), i
+
+ !$acc parallel loop present (data(m:m))
+ do i = m, m+m-1
+ data(i) = m
+ end do
+ !$acc end parallel loop
+end subroutine test_data
--
2.17.1
