Hi! On Fri, 30 Nov 2018 17:55:17 +0800, Chung-Lin Tang <chunglin_t...@mentor.com> wrote: > On 2018/7/21 6:07 AM, Cesar Philippidis wrote: > > This is another old gomp4 patch
The patch as posted here contains additional changes compared to the original patch. I removed these additional changes. The Fortran test case included doesn't actually FAIL without the patch, so I removed that one, too. The C/C++ test case included is enough to motivate the bug, and its fix. > > 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 wish PRs would get filed right when such things are discovered... I just filed <https://gcc.gnu.org/PR88288> "[OpenACC, libgomp] Adjust offsets for present data clauses". (Better late than never...) > 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. Thanks for having a look, I agree. This will likely need to be fixed on all active release branches. For now, committed to trunk in r266688: commit 2110057d427ba710d7f60045fe33c161e6b181c4 Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4> Date: Fri Nov 30 20:39:49 2018 +0000 [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 <tho...@codesourcery.com> git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@266688 138bc75d-0d04-0410-961f-82ee72b054a4 --- libgomp/ChangeLog | 6 ++++ libgomp/oacc-parallel.c | 3 +- .../testsuite/libgomp.oacc-c-c++-common/pr88288.c | 41 ++++++++++++++++++++++ 3 files changed, 49 insertions(+), 1 deletion(-) diff --git libgomp/ChangeLog libgomp/ChangeLog index a9dcbd808200..d095a196fb6c 100644 --- libgomp/ChangeLog +++ libgomp/ChangeLog @@ -1,3 +1,9 @@ +2018-11-30 Cesar Philippidis <ce...@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 <tho...@codesourcery.com> * testsuite/libgomp.oacc-fortran/lib-16-2.f90: New file. diff --git libgomp/oacc-parallel.c libgomp/oacc-parallel.c index b80ace585907..1e08af70b4da 100644 --- libgomp/oacc-parallel.c +++ 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 libgomp/testsuite/libgomp.oacc-c-c++-common/pr88288.c libgomp/testsuite/libgomp.oacc-c-c++-common/pr88288.c new file mode 100644 index 000000000000..d13e3359a3ec --- /dev/null +++ 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; +} Grüße Thomas
