Hi Julian!
On 2020-07-16T22:21:43+0100, Julian Brown <jul...@codesourcery.com> wrote:
> On Thu, 16 Jul 2020 11:35:23 +0200
> Thomas Schwinge <tho...@codesourcery.com> wrote:
>> 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)
Thanks for verifying.
> but I think in practice it's
> harmless (at least I haven't figured out a breaking test case).
I have. ;-P
(That's in a tree with the pending "[OpenACC] Deep copy attach/detach
should not affect reference counts" included -- may or may not be
relevant.)
> 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.
With the attached incremental patch merged in, OK for master and
releases/gcc-10 (once un-frozen) branches.
> --- 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);
ACK.
Regarding the 'k->refcount' that was used above...
> --- /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])
..., here we have 'k->refcount == 1', but...
> +
> + #pragma acc data copyin(mystr.arr[0:16])
> + {
..., here we now have 'k->refcount == 2' as the 'copyin' has incremented
it, so...
> + #pragma acc exit data detach(mystr.arr)
> + /* { dg-output "libgomp: attach count underflow" } */
> + }
..., it won't trigger the erroneous behavior here.
Instead of 'copyin(mystr.arr[0:16])' on the OpenACC 'data' construct, we
have to do a 'attach(mystr.arr)', and can then reproduce the problem:
without the 'libgomp/target.c:gomp_unmap_vars_internal' change, it
unexpectedly doesn't catch 'libgomp: attach count underflow', and instead
'detach'es again, and probably (potentially?) writes wrong data into
'mystr.arr' (I haven't tested that aspect).
> +
> + #pragma acc exit data copyout(mystr, localarr[0:16])
> +
> + return 0;
> +}
> +
> +/* { dg-shouldfail "" } */
Grüße
Thomas
-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander
Walter
>From 6f7f6f0ac80cdc96bc02777542297b6fef538c0b Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <tho...@codesourcery.com>
Date: Fri, 17 Jul 2020 09:41:18 +0200
Subject: [PATCH] into "openacc: Remove unnecessary detach finalization"
---
.../structured-detach-underflow.c | 12 ++++++++----
1 file changed, 8 insertions(+), 4 deletions(-)
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
index 184410137d8..fc1f59e2185 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-detach-underflow.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-detach-underflow.c
@@ -1,5 +1,7 @@
/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
+#include <stdio.h>
+
int main ()
{
struct {
@@ -10,15 +12,17 @@ int main ()
#pragma acc enter data copyin(mystr, localarr[0:16])
- #pragma acc data copyin(mystr.arr[0:16])
+ #pragma acc data attach(mystr.arr)
{
#pragma acc exit data detach(mystr.arr)
- /* { dg-output "libgomp: attach count underflow" } */
+ fprintf (stderr, "CheCKpOInT1\n");
+ /* { dg-output ".*CheCKpOInT1(\n|\r\n|\r)" } */
}
+ /* { dg-shouldfail "" }
+ { dg-output "(\n|\r\n|\r)libgomp: attach count underflow(\n|\r\n|\r)$" } */
+ fprintf (stderr, "CheCKpOInT2\n");
#pragma acc exit data copyout(mystr, localarr[0:16])
return 0;
}
-
-/* { dg-shouldfail "" } */
--
2.17.1
