Adjust offsets for present data clauses

Message ID b296bf2d-1071-93a7-17d7-c3d685bee220@codesourcery.com
State New
Headers show
Series
  • Adjust offsets for present data clauses
Related show

Commit Message

Cesar Philippidis July 20, 2018, 10:07 p.m.
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

Comments

Chung-Lin Tang Nov. 30, 2018, 9:55 a.m. | #1
On 2018/7/21 6:07 AM, Cesar Philippidis wrote:
> 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

> 


Hi Thomas, this patch should be within your maintainership area now.

I think this patch is pretty obvious; this is what the 'offset' field
of struct target_var_desc is supposed to be used for, and is in line
with other sites throughout target.c.

I do think it might be better to use a more succinct form like as attached,
you may consider which form better suits your taste when you apply it.

Chung-Lin
Index: libgomp/oacc-parallel.c
===================================================================
--- libgomp/oacc-parallel.c	(revision 266611)
+++ libgomp/oacc-parallel.c	(working copy)
@@ -231,8 +231,11 @@ 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);
+    devaddrs[i] = (tgt->list[i].key
+		   ? (void *) (tgt->list[i].key->tgt->tgt_start
+			       + tgt->list[i].key->tgt_offset
+			       + tgt->list[i].offset)
+		   : NULL);
 
   acc_dev->openacc.exec_func (tgt_fn, mapnum, hostaddrs, devaddrs,
 			      async, dims, tgt);
Julian Brown Nov. 30, 2018, 10:37 a.m. | #2
On Fri, 30 Nov 2018 17:55:17 +0800
Chung-Lin Tang <chunglin_tang@mentor.com> wrote:

> On 2018/7/21 6:07 AM, Cesar Philippidis wrote:

> > 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

> >   

> 

> Hi Thomas, this patch should be within your maintainership area now.

> 

> I think this patch is pretty obvious; this is what the 'offset' field

> of struct target_var_desc is supposed to be used for, and is in line

> with other sites throughout target.c.

> 

> I do think it might be better to use a more succinct form like as

> attached, you may consider which form better suits your taste when

> you apply it.


This one will be superseded by the attach/detach patch, I think (where
the additional offset is added also, via calling "gomp_map_val".

HTH,

Julian
Thomas Schwinge Nov. 30, 2018, 8:23 p.m. | #3
Hi Julian!

On Fri, 30 Nov 2018 10:37:09 +0000, Julian Brown <julian@codesourcery.com> wrote:
> > On 2018/7/21 6:07 AM, Cesar Philippidis wrote:

> > > 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>


> This one will be superseded


Well, this is clearly a bug, so should be addressed independently.  I
just filed <https://gcc.gnu.org/PR88288> "[OpenACC, libgomp] Adjust
offsets for present data clauses".  (Better late than never...)

> by the attach/detach patch, I think (where

> the additional offset is added also, via calling "gomp_map_val".


Aha, good to know: during review, I wondered about whether that function
should actually be used.  (But it also isn't in other places in
"libgomp/target.c".)


Grüße
 Thomas

Patch

From fb743d8a45193c177cb0082400d140949e8c1e6d Mon Sep 17 00:00:00 2001
From: Cesar Philippidis <cesar@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  <cesar@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