Hi Julian!
On 2020-06-16T15:38:33-0700, Julian Brown <jul...@codesourcery.com> wrote:
> This is a new version of the patch last sent here:
>
> https://gcc.gnu.org/pipermail/gcc-patches/2020-May/546332.html
>
> Minus the bits that Thomas has committed already (thanks!), and with
> adjustments to allow for GOMP_MAP_ATTACH being grouped together with a
> preceding clause.
>
> OK?
Please also update the "virtual refcount" comment in
'libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-4.c'.
Your patch now makes the 'libgomp.oacc-fortran/mdc-refcount-1-1-1.f90',
'libgomp.oacc-fortran/mdc-refcount-1-2-1.f90',
'libgomp.oacc-fortran/mdc-refcount-1-2-2.f90',
'libgomp.oacc-fortran/mdc-refcount-1-3-1.f90' test cases PASS (did you
not see that?), so we have to remove all XFAILing, 'print'/'dg-output'
etc. from these, and it changes the error reporting in
'libgomp.oacc-fortran/mdc-refcount-1-4-1.f90', so we have to adjust that.
See attached patch "into Adjust dynamic reference count semantics".
Basically OK for master branch and releases/gcc-10 branch. However,
still a few questions, which can be addressed first, or separately, as
appropriate:
> --- a/libgomp/oacc-mem.c
> +++ b/libgomp/oacc-mem.c
> @@ -1048,13 +1052,111 @@ goacc_enter_data_internal (struct gomp_device_descr
> *acc_dev, size_t mapnum,
> {
> for (size_t i = 0; i < mapnum; i++)
> {
> - int group_last = find_group_last (i, mapnum, sizes, kinds);
> + splay_tree_key n;
> + size_t group_last = find_group_last (i, mapnum, sizes, kinds);
> + bool struct_p = false;
> + size_t size, groupnum = (group_last - i) + 1;
> +
> + switch (kinds[i] & 0xff)
> + {
> + case GOMP_MAP_STRUCT:
> + {
> + int last = i + sizes[i];
(If you'd like to, see my comment about 'last' in
<87k10o72dd.fsf@euler.schwinge.homeip.net">http://mid.mail-archive.com/87k10o72dd.fsf@euler.schwinge.homeip.net>.)
> + size = (uintptr_t) hostaddrs[last] + sizes[last]
> + - (uintptr_t) hostaddrs[i];
> + struct_p = true;
> + }
> + break;
> +
> + case GOMP_MAP_ATTACH:
> + size = sizeof (void *);
> + break;
> +
> + default:
> + size = sizes[i];
> + }
> +
> + n = lookup_host (acc_dev, hostaddrs[i], size);
> +
> + if (n && struct_p)
> + {
> + if (n->refcount != REFCOUNT_INFINITY)
> + n->refcount += groupnum - 1;
> + n->dynamic_refcount += groupnum - 1;
> + gomp_mutex_unlock (&acc_dev->lock);
> + }
As that had already confused me before,
<87k10o72dd.fsf@euler.schwinge.homeip.net">http://mid.mail-archive.com/87k10o72dd.fsf@euler.schwinge.homeip.net>,
please add a minimal comment here, something like: "Increment refcount
not by one but by number of items in 'GOMP_MAP_STRUCT'".
> + else if (n && groupnum == 1)
> + {
> + void *h = hostaddrs[i];
> + size_t s = sizes[i];
> +
> + /* A standalone attach clause. */
> + if ((kinds[i] & 0xff) == GOMP_MAP_ATTACH)
> + gomp_attach_pointer (acc_dev, aq, &acc_dev->mem_map, n,
> + (uintptr_t) h, s, NULL);
> + else if (h + s > (void *) n->host_end)
> + {
> + gomp_mutex_unlock (&acc_dev->lock);
> + gomp_fatal ("[%p,+%d] not mapped", (void *)h, (int)s);
> + }
> +
> + assert (n->refcount != REFCOUNT_LINK);
> + if (n->refcount != REFCOUNT_INFINITY)
> + n->refcount++;
> + n->dynamic_refcount++;
>
> - gomp_map_vars_async (acc_dev, aq,
> - (group_last - i) + 1,
> - &hostaddrs[i], NULL,
> - &sizes[i], &kinds[i], true,
> - GOMP_MAP_VARS_OPENACC_ENTER_DATA);
> + gomp_mutex_unlock (&acc_dev->lock);
> + }
> + else if (n && groupnum > 1)
> + {
> + assert (n->refcount != REFCOUNT_INFINITY
> + && n->refcount != REFCOUNT_LINK);
> +
> + for (size_t j = i + 1; j <= group_last; j++)
> + if ((kinds[j] & 0xff) == GOMP_MAP_ATTACH)
> + {
> + splay_tree_key m
> + = lookup_host (acc_dev, hostaddrs[j], sizeof (void *));
> + gomp_attach_pointer (acc_dev, aq, &acc_dev->mem_map, m,
> + (uintptr_t) hostaddrs[j], sizes[j], NULL);
> + }
Per the earlier '[OpenACC] GOMP_MAP_ATTACH handling in find_group_last',
we should never have more than one 'GOMP_MAP_ATTACH' following something
else (right?), but it's still OK to leave this in this generic form --
unless you want to add some 'assert'ing here.
> +
> + bool processed = false;
> +
> + struct target_mem_desc *tgt = n->tgt;
> + for (size_t j = 0; j < tgt->list_count; j++)
> + if (tgt->list[j].key == n)
> + {
> + for (size_t k = 0; k < groupnum; k++)
> + if (j + k < tgt->list_count && tgt->list[j + k].key)
> + {
> + tgt->list[j + k].key->refcount++;
> + tgt->list[j + k].key->dynamic_refcount++;
> + }
> + processed = true;
> + }
> +
> + gomp_mutex_unlock (&acc_dev->lock);
> + if (!processed)
> + gomp_fatal ("dynamic refcount incrementing failed for "
> + "pointer/pset");
> + }
In <87k10o72dd.fsf@euler.schwinge.homeip.net">http://mid.mail-archive.com/87k10o72dd.fsf@euler.schwinge.homeip.net>
I had asked to "Please add some text to explain [...]" etc.
> + else if (hostaddrs[i])
> + {
> + gomp_mutex_unlock (&acc_dev->lock);
> +
> + struct target_mem_desc *tgt
> + = gomp_map_vars_async (acc_dev, aq, groupnum, &hostaddrs[i], NULL,
> + &sizes[i], &kinds[i], true,
> + GOMP_MAP_VARS_ENTER_DATA);
> + assert (tgt);
> + for (size_t j = 0; j < tgt->list_count; j++)
> + {
> + n = tgt->list[j].key;
> + if (n)
> + n->dynamic_refcount++;
> + }
> + }
In <87k10o72dd.fsf@euler.schwinge.homeip.net">http://mid.mail-archive.com/87k10o72dd.fsf@euler.schwinge.homeip.net>
I has asked to make this "else nothing" case more explicit -- if that's
correct, after all.
>
> i = group_last;
> }
Your patch regresses the attached
'libgomp.oacc-c-c++-common/struct-3-1-1.c', which used to act like
detailed in the file, but now does:
CheCKpOInT1
CheCKpOInT2
a.out:
source-gcc/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-3-1-1.c:28: main:
Assertion `acc_is_present (&s.b, sizeof s.b)' failed.
Aborted (core dumped)
That means, after '#pragma acc enter data create(s.a)' we're no longer
refusing '#pragma acc enter data create(s.b)', but then the
'acc_is_present' for 's.b' fails. Is that a true regression introduced
by your patch, or a separate issue (which before just worked by chance)?
In the latter case, please file a PR.
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 31f7b23a8ec4107898d612f2c758f39faa0f0691 Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <tho...@codesourcery.com>
Date: Tue, 30 Jun 2020 15:48:37 +0200
Subject: [PATCH] into Adjust dynamic reference count semantics
---
.../testsuite/libgomp.oacc-fortran/mdc-refcount-1-1-1.f90 | 8 --------
.../testsuite/libgomp.oacc-fortran/mdc-refcount-1-1-2.F90 | 5 +----
.../testsuite/libgomp.oacc-fortran/mdc-refcount-1-2-1.f90 | 8 --------
.../testsuite/libgomp.oacc-fortran/mdc-refcount-1-2-2.f90 | 8 --------
.../testsuite/libgomp.oacc-fortran/mdc-refcount-1-3-1.f90 | 8 --------
.../testsuite/libgomp.oacc-fortran/mdc-refcount-1-4-1.f90 | 7 +++----
6 files changed, 4 insertions(+), 40 deletions(-)
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-1-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-1-1.f90
index 445cbabb8ca..1d97dd382d4 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-1-1.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-1-1.f90
@@ -21,15 +21,7 @@ program main
if (.not. acc_is_present(var%a)) stop 1
if (.not. acc_is_present(var)) stop 2
- print *, "CheCKpOInT1"
- ! { dg-output ".*CheCKpOInT1(\n|\r\n|\r)" }
!$acc exit data delete(var%a) finalize
- !TODO goacc_exit_data_internal: Assertion `is_tgt_unmapped || num_mappings > 1' failed.
- !TODO { dg-output ".*\[Aa\]ssert.*is_tgt_unmapped" { target { ! openacc_host_selected } } } ! Scan for what we expect in the "XFAILed" case (without actually XFAILing).
- !TODO { dg-shouldfail "XFAILed" { ! openacc_host_selected } } ! ... instead of 'dg-xfail-run-if' so that 'dg-output' is evaluated at all.
- !TODO { dg-final { if { [dg-process-target { xfail { ! openacc_host_selected } }] == "F" } { xfail "[testname-for-summary] really is XFAILed" } } } ! ... so that we still get an XFAIL visible in the log.
- print *, "CheCKpOInT2"
- ! { dg-output ".CheCKpOInT2(\n|\r\n|\r)" { target { openacc_host_selected } } }
if (acc_is_present(var%a)) stop 3
if (.not. acc_is_present(var)) stop 4
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-1-2.F90 b/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-1-2.F90
index 7b206ac2042..4307f50c46e 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-1-2.F90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-1-2.F90
@@ -1,9 +1,6 @@
! { dg-do run }
! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } }
-/* Nullify the 'finalize' clause, which disturbs reference counting. */
+/* Nullify the 'finalize' clause. */
#define finalize
#include "mdc-refcount-1-1-1.f90"
-
-! { dg-output ".*CheCKpOInT1(\n|\r\n|\r)" }
-! { dg-output ".CheCKpOInT2(\n|\r\n|\r)" }
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-2-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-2-1.f90
index 8554534b2f2..e6f3f4afc3b 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-2-1.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-2-1.f90
@@ -23,15 +23,7 @@ program main
if (.not. acc_is_present(var%a)) stop 1
if (.not. acc_is_present(var)) stop 2
- print *, "CheCKpOInT1"
- ! { dg-output ".*CheCKpOInT1(\n|\r\n|\r)" }
!$acc exit data delete(var%a) finalize
- !TODO goacc_exit_data_internal: Assertion `is_tgt_unmapped || num_mappings > 1' failed.
- !TODO { dg-output ".*\[Aa\]ssert.*is_tgt_unmapped" { target { ! openacc_host_selected } } } ! Scan for what we expect in the "XFAILed" case (without actually XFAILing).
- !TODO { dg-shouldfail "XFAILed" { ! openacc_host_selected } } ! ... instead of 'dg-xfail-run-if' so that 'dg-output' is evaluated at all.
- !TODO { dg-final { if { [dg-process-target { xfail { ! openacc_host_selected } }] == "F" } { xfail "[testname-for-summary] really is XFAILed" } } } ! ... so that we still get an XFAIL visible in the log.
- print *, "CheCKpOInT2"
- ! { dg-output ".CheCKpOInT2(\n|\r\n|\r)" { target { openacc_host_selected } } }
if (acc_is_present(var%a)) stop 3
if (.not. acc_is_present(var)) stop 4
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-2-2.f90 b/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-2-2.f90
index 8e696cc70e8..78f54e64dce 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-2-2.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-2-2.f90
@@ -23,15 +23,7 @@ program main
if (.not. acc_is_present(var%a)) stop 1
if (.not. acc_is_present(var)) stop 2
- print *, "CheCKpOInT1"
- ! { dg-output ".*CheCKpOInT1(\n|\r\n|\r)" }
!$acc exit data delete(var%a)
- !TODO goacc_exit_data_internal: Assertion `is_tgt_unmapped || num_mappings > 1' failed.
- !TODO { dg-output ".*\[Aa\]ssert.*is_tgt_unmapped" { target { ! openacc_host_selected } } } ! Scan for what we expect in the "XFAILed" case (without actually XFAILing).
- !TODO { dg-shouldfail "XFAILed" { ! openacc_host_selected } } ! ... instead of 'dg-xfail-run-if' so that 'dg-output' is evaluated at all.
- !TODO { dg-final { if { [dg-process-target { xfail { ! openacc_host_selected } }] == "F" } { xfail "[testname-for-summary] really is XFAILed" } } } ! ... so that we still get an XFAIL visible in the log.
- print *, "CheCKpOInT2"
- ! { dg-output ".CheCKpOInT2(\n|\r\n|\r)" { target { openacc_host_selected } } }
if (acc_is_present(var%a)) stop 3
if (.not. acc_is_present(var)) stop 4
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-3-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-3-1.f90
index 070a6f8e149..f9dcb485b26 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-3-1.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-3-1.f90
@@ -24,15 +24,7 @@ program main
if (.not. acc_is_present(var)) stop 2
!$acc exit data detach(var%a)
- print *, "CheCKpOInT1"
- ! { dg-output ".*CheCKpOInT1(\n|\r\n|\r)" }
!$acc exit data delete(var%a) finalize
- !TODO goacc_exit_data_internal: Assertion `is_tgt_unmapped || num_mappings > 1' failed.
- !TODO { dg-output ".*\[Aa\]ssert.*is_tgt_unmapped" { target { ! openacc_host_selected } } } ! Scan for what we expect in the "XFAILed" case (without actually XFAILing).
- !TODO { dg-shouldfail "XFAILed" { ! openacc_host_selected } } ! ... instead of 'dg-xfail-run-if' so that 'dg-output' is evaluated at all.
- !TODO { dg-final { if { [dg-process-target { xfail { ! openacc_host_selected } }] == "F" } { xfail "[testname-for-summary] really is XFAILed" } } } ! ... so that we still get an XFAIL visible in the log.
- print *, "CheCKpOInT2"
- ! { dg-output ".CheCKpOInT2(\n|\r\n|\r)" { target { openacc_host_selected } } }
if (acc_is_present(var%a)) stop 3
if (.not. acc_is_present(var)) stop 4
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-4-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-4-1.f90
index b22e411567f..fbd52373946 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-4-1.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-4-1.f90
@@ -23,16 +23,15 @@ program main
if (.not. acc_is_present(var%a)) stop 1
if (.not. acc_is_present(var)) stop 2
+ !$acc exit data detach(var%a) finalize
print *, "CheCKpOInT1"
! { dg-output ".*CheCKpOInT1(\n|\r\n|\r)" }
- !$acc exit data detach(var%a) finalize
- !TODO goacc_exit_data_internal: Assertion `is_tgt_unmapped || num_mappings > 1' failed.
- !TODO { dg-output ".*\[Aa\]ssert.*is_tgt_unmapped" { target { ! openacc_host_selected } } } ! Scan for what we expect in the "XFAILed" case (without actually XFAILing).
+ !$acc exit data delete(var%a)
+ !TODO { dg-output "(\n|\r\n|\r)libgomp: attach count underflow(\n|\r\n|\r)$" { target { ! openacc_host_selected } } } ! Scan for what we expect in the "XFAILed" case (without actually XFAILing).
!TODO { dg-shouldfail "XFAILed" { ! openacc_host_selected } } ! ... instead of 'dg-xfail-run-if' so that 'dg-output' is evaluated at all.
!TODO { dg-final { if { [dg-process-target { xfail { ! openacc_host_selected } }] == "F" } { xfail "[testname-for-summary] really is XFAILed" } } } ! ... so that we still get an XFAIL visible in the log.
print *, "CheCKpOInT2"
! { dg-output ".CheCKpOInT2(\n|\r\n|\r)" { target { openacc_host_selected } } }
- !$acc exit data delete(var%a)
if (acc_is_present(var%a)) stop 3
if (.not. acc_is_present(var)) stop 4
--
2.17.1
/* Test dynamic mapping of separate structure members. */
#include <assert.h>
#include <stdio.h>
#include <openacc.h>
struct s
{
char a;
float b;
};
int main ()
{
struct s s;
#pragma acc enter data create(s.a)
assert (acc_is_present (&s.a, sizeof s.a));
fprintf (stderr, "CheCKpOInT1\n");
/* { dg-output ".*CheCKpOInT1(\n|\r\n|\r)" } */
#pragma acc enter data create(s.b)
/* { dg-output "(\n|\r\n|\r)libgomp: Trying to map into device \\\[\[0-9a-fA-FxX.\]+\\\) structure element when other mapped elements from the same structure weren't mapped together with it(\n|\r\n|\r)$" { target { ! openacc_host_selected } } } ! Scan for what we expect in the "XFAILed" case (without actually XFAILing).
{ dg-shouldfail "XFAILed" { ! openacc_host_selected } } ! ... instead of 'dg-xfail-run-if' so that 'dg-output' is evaluated at all.
{ dg-final { if { [dg-process-target { xfail { ! openacc_host_selected } }] == "F" } { xfail "[testname-for-summary] really is XFAILed" } } } ! ... so that we still get an XFAIL visible in the log. */
fprintf (stderr, "CheCKpOInT2\n");
/* { dg-output "CheCKpOInT2(\n|\r\n|\r)" { target { openacc_host_selected } } } */
assert (acc_is_present (&s.b, sizeof s.b));
//TODO PR95236
assert (acc_is_present (&s, sizeof s));
return 0;
}
