On Thu, 16 Jul 2020 11:35:23 +0200
Thomas Schwinge <tho...@codesourcery.com> wrote:
> Hi Julian!
>
> Ping.
>
> On 2020-06-26T11:20:40+0200, I wrote:
> > On 2019-12-17T22:03:47-0800, Julian Brown <jul...@codesourcery.com>
> > wrote:
> >> This part contains the libgomp runtime support for the
> >> GOMP_MAP_ATTACH and GOMP_MAP_DETACH mapping kinds (etc.), as
> >> introduced by the front-end patches following in this series.
> >
> >> --- a/libgomp/target.c
> >> +++ b/libgomp/target.c
> >
> >> @@ -1534,6 +1571,18 @@ gomp_unmap_vars_internal (struct
> >> target_mem_desc *tgt, bool do_copyfrom,
> >
> > This is the code path at the end of a structured OpenACC 'data'
> > construct.
> >
> >> + /* We must perform detachments before any copies back to the
> >> host. */
> >> + for (i = 0; i < tgt->list_count; i++)
> >> + {
> >> + splay_tree_key k = tgt->list[i].key;
> >> +
> >> + if (k != NULL && tgt->list[i].do_detach)
> >> + gomp_detach_pointer (devicep, aq, k,
> >> tgt->list[i].key->host_start
> >> + +
> >> tgt->list[i].offset,
> >> + k->refcount == 1, NULL);
> >> + }
> >
> > Can you please explain (as a source code comment) the logic for here
> > using 'k->refcount == 1' for the 'bool finalize' parameter of
> > 'gomp_detach_pointer'; this somehow feels "strange"?
[snip]
> > Shouldn't this just always be 'finalize = false' given that there
> > is no 'finalize' semantics for 'detach' on a structured OpenACC
> > 'data' constructs -- at least that's what I remember right now?
As far as I can tell, forcing finalize there is unnecessary (and as you
point out, conceptually dubious), but I think in practice it's
harmless (at least I haven't figured out a breaking test case). Anyway,
this patch just passes "false" for the finalize argument. I've also
added a test case, though it passes before/after the patch.
OK? Tested with offloading to nvptx.
Julian
>From 31618a79bb3aa0d088030904fff0ad386ddb0999 Mon Sep 17 00:00:00 2001
From: Julian Brown <jul...@codesourcery.com>
Date: Thu, 2 Jul 2020 14:18:20 -0700
Subject: [PATCH] openacc: Remove unnecessary detach finalization
The call to gomp_detach_pointer in gomp_unmap_vars_internal does not
need to force finalization, and doing so may mask mismatched pointer
attachments/detachments. This patch removes the forcing.
2020-07-16 Julian Brown <jul...@codesourcery.com>
libgomp/
* target.c (gomp_unmap_vars_internal): Remove unnecessary forcing of
finalization for detach operation.
* testsuite/libgomp.oacc-c-c++-common/structured-detach-underflow.c:
New test.
---
libgomp/target.c | 2 +-
.../structured-detach-underflow.c | 24 +++++++++++++++++++
2 files changed, 25 insertions(+), 1 deletion(-)
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/structured-detach-underflow.c
diff --git a/libgomp/target.c b/libgomp/target.c
index d6b3572c8d8..00c75fbd885 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -1437,7 +1437,7 @@ gomp_unmap_vars_internal (struct target_mem_desc *tgt, bool do_copyfrom,
if (k != NULL && tgt->list[i].do_detach)
gomp_detach_pointer (devicep, aq, k, tgt->list[i].key->host_start
+ tgt->list[i].offset,
- k->refcount == 1, NULL);
+ false, NULL);
}
for (i = 0; i < tgt->list_count; i++)
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-detach-underflow.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-detach-underflow.c
new file mode 100644
index 00000000000..184410137d8
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-detach-underflow.c
@@ -0,0 +1,24 @@
+/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
+
+int main ()
+{
+ struct {
+ int *arr;
+ } mystr;
+ int localarr[16];
+ mystr.arr = localarr;
+
+ #pragma acc enter data copyin(mystr, localarr[0:16])
+
+ #pragma acc data copyin(mystr.arr[0:16])
+ {
+ #pragma acc exit data detach(mystr.arr)
+ /* { dg-output "libgomp: attach count underflow" } */
+ }
+
+ #pragma acc exit data copyout(mystr, localarr[0:16])
+
+ return 0;
+}
+
+/* { dg-shouldfail "" } */
--
2.23.0
