From: Thomas Schwinge <thomas@codesourcery.com>
To: Julian Brown <julian@codesourcery.com>, <gcc-patches@gcc.gnu.org>
Cc: Jakub Jelinek <jakub@redhat.com>
Subject: [OpenACC] Repair/restore 'is_tgt_unmapped' checking (was: [PATCH 07/13] OpenACC 2.6 deep copy: libgomp parts)
Date: Thu, 4 Jun 2020 20:35:17 +0200 [thread overview]
Message-ID: <87bllysmqi.fsf@euler.schwinge.homeip.net> (raw)
In-Reply-To: <20200520201100.0d5f1f2f@squid.athome>
[-- Attachment #1: Type: text/plain, Size: 4251 bytes --]
Hi!
On 2020-05-20T20:11:00+0100, Julian Brown <julian@codesourcery.com> wrote:
> On Wed, 20 May 2020 16:52:02 +0200
> Thomas Schwinge <thomas@codesourcery.com> wrote:
>> On 2019-12-17T22:03:47-0800, Julian Brown <julian@codesourcery.com>
>> wrote:
>> > --- a/libgomp/oacc-mem.c
>> > +++ b/libgomp/oacc-mem.c
>>
>> > static int
>> > -find_group_last (int pos, size_t mapnum, unsigned short *kinds)
>> > +find_group_last (int pos, size_t mapnum, size_t *sizes, unsigned short *kinds) {
>> > unsigned char kind0 = kinds[pos] & 0xff;
>> > - int first_pos = pos, last_pos = pos;
>> > + int first_pos = pos;
>> >
>> > - if (kind0 == GOMP_MAP_TO_PSET)
>> > + switch (kind0)
>> > {
>> > + case GOMP_MAP_TO_PSET:
>> > while (pos + 1 < mapnum && (kinds[pos + 1] & 0xff) == GOMP_MAP_POINTER)
>> > - last_pos = ++pos;
>> > + pos++;
>> > /* We expect at least one GOMP_MAP_POINTER after a GOMP_MAP_TO_PSET. */
>> > - assert (last_pos > first_pos);
>> > - }
>> > - else
>> > - {
>> > + assert (pos > first_pos);
>> > + break;
>> > +
>> > + case GOMP_MAP_STRUCT:
>> > + pos += sizes[pos];
>> > + break;
>> > +
>> > + case GOMP_MAP_POINTER:
>> > + case GOMP_MAP_ALWAYS_POINTER:
>> > + /* These mappings are only expected after some other mapping. If we
>> > + see one by itself, something has gone wrong. */
>> > + gomp_fatal ("unexpected mapping");
>> > + break;
>> > +
>> > + default:
>> > /* GOMP_MAP_ALWAYS_POINTER can only appear directly after some other mapping. */
>> > - if (pos + 1 < mapnum
>> > - && (kinds[pos + 1] & 0xff) == GOMP_MAP_ALWAYS_POINTER)
>> > - return pos + 1;
>> > + if (pos + 1 < mapnum)
>> > + {
>> > + unsigned char kind1 = kinds[pos + 1] & 0xff;
>> > + if (kind1 == GOMP_MAP_ALWAYS_POINTER)
>> > + return pos + 1;
>> > + }
>> >
>> > - /* We can have one or several GOMP_MAP_POINTER mappings after a to/from
>> > + /* We can have zero or more GOMP_MAP_POINTER mappings after a to/from (etc.) mapping. */
>> > while (pos + 1 < mapnum && (kinds[pos + 1] & 0xff) == GOMP_MAP_POINTER)
>> > - last_pos = ++pos;
>> > + pos++;
>> > }
>> >
>> > - return last_pos;
>> > + return pos;
>> > }
>>
>> So this now causes grouped (!) mapping of all of 'GOMP_MAP_STRUCT',
>> that is, all its "members" at once.
>>
>> This, I suppose, mandated the removal of (some of) the
>> 'is_tgt_unmapped' checking (unfortunately committed not here, but as
>> part of r279621 "OpenACC reference count overhaul"), where we had
>> unmapping code (conceptually) similar to:
>>
>> bool is_tgt_unmapped = gomp_remove_var (acc_dev, n);
>> assert (is_tgt_unmapped);
>>
>> I'd introduced this a little bit earlier, finding this a simple yet
>> effective run-time, low-overhead consistency checking of (certain
>> aspects of) reference counting -- so just noting here that it's
>> somewhat bad that we can't have this anymore "just" because of
>> 'GOMP_MAP_STRUCT'. (Maybe there is a way to get it back; that's for
>> later?)
>
> I'm actually looking at this now as part of revisiting the refcounting
> work. I'm seeing what I can come up with in terms of being able to keep
> the runtime test (and fixing the other part you mentioned).
Good idea to skip the checking if 'num_mappings > 1', thanks! You've
included these changes as part of a different, bigger patch; I've split
them out, and pushed "[OpenACC] Repair/restore 'is_tgt_unmapped'
checking" to master branch in commit
06ec61726d192659cd446e59a91e78745037f0fd, and releases/gcc-10 branch in
commit 125621f569cfac9f4caa6afc1976d42b3d21359e, see attached.
Also note how this checking triggers for OpenACC 'finalize' clause usage
in 'libgomp.oacc-fortran/deep-copy-6.f90' (which I had relatedly XFAILed
before); so good to see that it's useful for something! (..., and did
your reference count self-checking not catch this case?)
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
[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #2: 0001-OpenACC-Repair-restore-is_tgt_unmapped-checking.patch --]
[-- Type: text/x-diff, Size: 16571 bytes --]
From 06ec61726d192659cd446e59a91e78745037f0fd Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <thomas@codesourcery.com>
Date: Fri, 29 May 2020 15:22:42 +0200
Subject: [PATCH] [OpenACC] Repair/restore 'is_tgt_unmapped' checking
libgomp/
* oacc-mem.c (goacc_exit_datum): Repair 'is_tgt_unmapped'
checking.
(acc_unmap_data, goacc_exit_data_internal): Restore
'is_tgt_unmapped' checking.
* testsuite/libgomp.oacc-c-c++-common/struct-refcount-1.c: New
file.
* testsuite/libgomp.oacc-fortran/deep-copy-6.f90: Adjust.
* testsuite/libgomp.oacc-fortran/mdc-refcount-1-1-1.f90: Likewise.
* testsuite/libgomp.oacc-fortran/mdc-refcount-1-2-1.f90: Likewise.
* testsuite/libgomp.oacc-fortran/mdc-refcount-1-2-2.f90: Likewise.
* testsuite/libgomp.oacc-fortran/mdc-refcount-1-3-1.f90: Likewise.
* testsuite/libgomp.oacc-fortran/mdc-refcount-1-4-1.f90: Likewise.
Co-Authored-By: Julian Brown <julian@codesourcery.com>
---
libgomp/oacc-mem.c | 60 +++++++++++++++++--
.../struct-refcount-1.c | 47 +++++++++++++++
.../libgomp.oacc-fortran/deep-copy-6.f90 | 15 ++---
.../mdc-refcount-1-1-1.f90 | 9 +--
.../mdc-refcount-1-2-1.f90 | 9 +--
.../mdc-refcount-1-2-2.f90 | 9 +--
.../mdc-refcount-1-3-1.f90 | 9 +--
.../mdc-refcount-1-4-1.f90 | 7 ++-
8 files changed, 135 insertions(+), 30 deletions(-)
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/struct-refcount-1.c
diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c
index 8e8c7c3093d5..b7c85cf5976f 100644
--- a/libgomp/oacc-mem.c
+++ b/libgomp/oacc-mem.c
@@ -485,7 +485,8 @@ acc_unmap_data (void *h)
tgt->tgt_end = 0;
tgt->to_free = NULL;
- gomp_remove_var (acc_dev, n);
+ bool is_tgt_unmapped = gomp_remove_var (acc_dev, n);
+ assert (is_tgt_unmapped);
gomp_mutex_unlock (&acc_dev->lock);
@@ -727,8 +728,16 @@ goacc_exit_datum (void *h, size_t s, unsigned short kind, int async)
gomp_remove_var_async (acc_dev, n, aq);
else
{
+ size_t num_mappings = 0;
+ /* If the target_mem_desc represents a single data mapping, we can
+ check that it is freed when this splay tree key's refcount reaches
+ zero. Otherwise (e.g. for a 'GOMP_MAP_STRUCT' mapping with
+ multiple members), fall back to skipping the test. */
+ for (size_t l_i = 0; l_i < n->tgt->list_count; ++l_i)
+ if (n->tgt->list[l_i].key)
+ ++num_mappings;
bool is_tgt_unmapped = gomp_remove_var (acc_dev, n);
- assert (is_tgt_unmapped);
+ assert (is_tgt_unmapped || num_mappings > 1);
}
}
@@ -1145,7 +1154,28 @@ goacc_exit_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum,
cur_node.host_end - cur_node.host_start);
if (n->refcount == 0)
- gomp_remove_var_async (acc_dev, n, aq);
+ {
+ if (aq)
+ /* TODO We can't do the 'is_tgt_unmapped' checking -- see the
+ 'gomp_unref_tgt' comment in
+ <http://mid.mail-archive.com/878snl36eu.fsf@euler.schwinge.homeip.net>;
+ PR92881. */
+ gomp_remove_var_async (acc_dev, n, aq);
+ else
+ {
+ size_t num_mappings = 0;
+ /* If the target_mem_desc represents a single data mapping,
+ we can check that it is freed when this splay tree key's
+ refcount reaches zero. Otherwise (e.g. for a
+ 'GOMP_MAP_STRUCT' mapping with multiple members), fall
+ back to skipping the test. */
+ for (size_t l_i = 0; l_i < n->tgt->list_count; ++l_i)
+ if (n->tgt->list[l_i].key)
+ ++num_mappings;
+ bool is_tgt_unmapped = gomp_remove_var (acc_dev, n);
+ assert (is_tgt_unmapped || num_mappings > 1);
+ }
+ }
}
break;
@@ -1177,7 +1207,29 @@ goacc_exit_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum,
&& str->refcount != REFCOUNT_INFINITY)
str->refcount--;
if (str->refcount == 0)
- gomp_remove_var_async (acc_dev, str, aq);
+ {
+ if (aq)
+ /* TODO We can't do the 'is_tgt_unmapped' checking --
+ see the 'gomp_unref_tgt' comment in
+ <http://mid.mail-archive.com/878snl36eu.fsf@euler.schwinge.homeip.net>;
+ PR92881. */
+ gomp_remove_var_async (acc_dev, str, aq);
+ else
+ {
+ size_t num_mappings = 0;
+ /* If the target_mem_desc represents a single data
+ mapping, we can check that it is freed when this
+ splay tree key's refcount reaches zero.
+ Otherwise (e.g. for a 'GOMP_MAP_STRUCT' mapping
+ with multiple members), fall back to skipping
+ the test. */
+ for (size_t l_i = 0; l_i < str->tgt->list_count; ++l_i)
+ if (str->tgt->list[l_i].key)
+ ++num_mappings;
+ bool is_tgt_unmapped = gomp_remove_var (acc_dev, str);
+ assert (is_tgt_unmapped || num_mappings > 1);
+ }
+ }
}
}
i += elems;
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-refcount-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-refcount-1.c
new file mode 100644
index 000000000000..bde5890d6676
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-refcount-1.c
@@ -0,0 +1,47 @@
+/* Test dynamic unmapping of separate structure members. */
+
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+
+#include <assert.h>
+#include <stdbool.h>
+#include <openacc.h>
+
+struct s
+{
+ char a;
+ float b;
+};
+
+void test (bool use_directives)
+{
+ struct s s;
+
+#pragma acc enter data create(s.a, s.b)
+ assert (acc_is_present (&s.a, sizeof s.a));
+ assert (acc_is_present (&s.b, sizeof s.b));
+
+ if (use_directives)
+ {
+#pragma acc exit data delete(s.a)
+ }
+ else
+ acc_delete (&s.a, sizeof s.a);
+ assert (!acc_is_present (&s.a, sizeof s.a));
+ assert (acc_is_present (&s.b, sizeof s.b));
+ if (use_directives)
+ {
+#pragma acc exit data delete(s.b)
+ }
+ else
+ acc_delete (&s.b, sizeof s.b);
+ assert (!acc_is_present (&s.a, sizeof s.a));
+ assert (!acc_is_present (&s.b, sizeof s.b));
+}
+
+int main ()
+{
+ test (true);
+ test (false);
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-6.f90 b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-6.f90
index a7943d93d542..5837a4039106 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-6.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-6.f90
@@ -40,19 +40,20 @@ program dtype
if (.not. acc_is_present(var%a(5:n - 5))) stop 11
if (.not. acc_is_present(var%b(5:n - 5))) stop 12
if (.not. acc_is_present(var)) stop 13
-!$acc exit data copyout(var%a(5:n - 5), var%b(5:n - 5)) finalize
- if (acc_get_device_type() .ne. acc_device_host) then
- if (acc_is_present(var%a(5:n - 5))) stop 21
- if (acc_is_present(var%b(5:n - 5))) stop 22
- end if
print *, "CheCKpOInT1"
! { dg-output ".*CheCKpOInT1(\n|\r\n|\r)" }
- if (.not. acc_is_present(var)) stop 23
- !TODO { dg-output "STOP 23(\n|\r\n|\r)$" { target { ! openacc_host_selected } } } ! Scan for what we expect in the "XFAILed" case (without actually XFAILing).
+!$acc exit data copyout(var%a(5:n - 5), var%b(5:n - 5)) 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_get_device_type() .ne. acc_device_host) then
+ if (acc_is_present(var%a(5:n - 5))) stop 21
+ if (acc_is_present(var%b(5:n - 5))) stop 22
+ end if
+ if (.not. acc_is_present(var)) stop 23
!$acc end data
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 449a6cf6894e..445cbabb8ca4 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,16 +21,17 @@ program main
if (.not. acc_is_present(var%a)) stop 1
if (.not. acc_is_present(var)) stop 2
- !$acc exit data delete(var%a) finalize
- if (acc_is_present(var%a)) stop 3
print *, "CheCKpOInT1"
! { dg-output ".*CheCKpOInT1(\n|\r\n|\r)" }
- if (.not. acc_is_present(var)) stop 4
- !TODO { dg-output "STOP 4(\n|\r\n|\r)$" { target { ! openacc_host_selected } } } ! Scan for what we expect in the "XFAILed" case (without actually XFAILing).
+ !$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
!$acc end data
if (acc_is_present(var%a)) stop 5
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 a7e649d30415..8554534b2f27 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,16 +23,17 @@ program main
if (.not. acc_is_present(var%a)) stop 1
if (.not. acc_is_present(var)) stop 2
- !$acc exit data delete(var%a) finalize
- if (acc_is_present(var%a)) stop 3
print *, "CheCKpOInT1"
! { dg-output ".*CheCKpOInT1(\n|\r\n|\r)" }
- if (.not. acc_is_present(var)) stop 4
- !TODO { dg-output "STOP 4(\n|\r\n|\r)$" { target { ! openacc_host_selected } } } ! Scan for what we expect in the "XFAILed" case (without actually XFAILing).
+ !$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
!$acc end data
if (acc_is_present(var%a)) stop 5
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 3402fafc7e28..8e696cc70e80 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,16 +23,17 @@ program main
if (.not. acc_is_present(var%a)) stop 1
if (.not. acc_is_present(var)) stop 2
- !$acc exit data delete(var%a)
- if (acc_is_present(var%a)) stop 3
print *, "CheCKpOInT1"
! { dg-output ".*CheCKpOInT1(\n|\r\n|\r)" }
- if (.not. acc_is_present(var)) stop 4
- !TODO { dg-output "STOP 4(\n|\r\n|\r)$" { target { ! openacc_host_selected } } } ! Scan for what we expect in the "XFAILed" case (without actually XFAILing).
+ !$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
!$acc end data
if (acc_is_present(var%a)) stop 5
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 7504969d9a56..070a6f8e1493 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,16 +24,17 @@ program main
if (.not. acc_is_present(var)) stop 2
!$acc exit data detach(var%a)
- !$acc exit data delete(var%a) finalize
- if (acc_is_present(var%a)) stop 3
print *, "CheCKpOInT1"
! { dg-output ".*CheCKpOInT1(\n|\r\n|\r)" }
- if (.not. acc_is_present(var)) stop 4
- !TODO { dg-output "STOP 4(\n|\r\n|\r)$" { target { ! openacc_host_selected } } } ! Scan for what we expect in the "XFAILed" case (without actually XFAILing).
+ !$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
!$acc end data
if (acc_is_present(var%a)) stop 5
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 fedae0db054b..b22e411567ff 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,15 +23,16 @@ 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 delete(var%a)
- !TODO { dg-output "(\n|\r\n|\r)libgomp: struct not mapped for detach operation(\n|\r\n|\r)$" { target { ! openacc_host_selected } } } ! Scan for what we expect in the "XFAILed" case (without actually XFAILing).
+ !$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).
!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.26.2
[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #3: 0001-OpenACC-Repair-restore-is_tgt_unmapped-checking.g10.patch --]
[-- Type: text/x-diff, Size: 16640 bytes --]
From 125621f569cfac9f4caa6afc1976d42b3d21359e Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <thomas@codesourcery.com>
Date: Fri, 29 May 2020 15:22:42 +0200
Subject: [PATCH] [OpenACC] Repair/restore 'is_tgt_unmapped' checking
libgomp/
* oacc-mem.c (goacc_exit_datum): Repair 'is_tgt_unmapped'
checking.
(acc_unmap_data, goacc_exit_data_internal): Restore
'is_tgt_unmapped' checking.
* testsuite/libgomp.oacc-c-c++-common/struct-refcount-1.c: New
file.
* testsuite/libgomp.oacc-fortran/deep-copy-6.f90: Adjust.
* testsuite/libgomp.oacc-fortran/mdc-refcount-1-1-1.f90: Likewise.
* testsuite/libgomp.oacc-fortran/mdc-refcount-1-2-1.f90: Likewise.
* testsuite/libgomp.oacc-fortran/mdc-refcount-1-2-2.f90: Likewise.
* testsuite/libgomp.oacc-fortran/mdc-refcount-1-3-1.f90: Likewise.
* testsuite/libgomp.oacc-fortran/mdc-refcount-1-4-1.f90: Likewise.
Co-Authored-By: Julian Brown <julian@codesourcery.com>
(cherry picked from commit 06ec61726d192659cd446e59a91e78745037f0fd)
---
libgomp/oacc-mem.c | 60 +++++++++++++++++--
.../struct-refcount-1.c | 47 +++++++++++++++
.../libgomp.oacc-fortran/deep-copy-6.f90 | 15 ++---
.../mdc-refcount-1-1-1.f90 | 9 +--
.../mdc-refcount-1-2-1.f90 | 9 +--
.../mdc-refcount-1-2-2.f90 | 9 +--
.../mdc-refcount-1-3-1.f90 | 9 +--
.../mdc-refcount-1-4-1.f90 | 7 ++-
8 files changed, 135 insertions(+), 30 deletions(-)
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/struct-refcount-1.c
diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c
index 8e8c7c3093d5..b7c85cf5976f 100644
--- a/libgomp/oacc-mem.c
+++ b/libgomp/oacc-mem.c
@@ -485,7 +485,8 @@ acc_unmap_data (void *h)
tgt->tgt_end = 0;
tgt->to_free = NULL;
- gomp_remove_var (acc_dev, n);
+ bool is_tgt_unmapped = gomp_remove_var (acc_dev, n);
+ assert (is_tgt_unmapped);
gomp_mutex_unlock (&acc_dev->lock);
@@ -727,8 +728,16 @@ goacc_exit_datum (void *h, size_t s, unsigned short kind, int async)
gomp_remove_var_async (acc_dev, n, aq);
else
{
+ size_t num_mappings = 0;
+ /* If the target_mem_desc represents a single data mapping, we can
+ check that it is freed when this splay tree key's refcount reaches
+ zero. Otherwise (e.g. for a 'GOMP_MAP_STRUCT' mapping with
+ multiple members), fall back to skipping the test. */
+ for (size_t l_i = 0; l_i < n->tgt->list_count; ++l_i)
+ if (n->tgt->list[l_i].key)
+ ++num_mappings;
bool is_tgt_unmapped = gomp_remove_var (acc_dev, n);
- assert (is_tgt_unmapped);
+ assert (is_tgt_unmapped || num_mappings > 1);
}
}
@@ -1145,7 +1154,28 @@ goacc_exit_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum,
cur_node.host_end - cur_node.host_start);
if (n->refcount == 0)
- gomp_remove_var_async (acc_dev, n, aq);
+ {
+ if (aq)
+ /* TODO We can't do the 'is_tgt_unmapped' checking -- see the
+ 'gomp_unref_tgt' comment in
+ <http://mid.mail-archive.com/878snl36eu.fsf@euler.schwinge.homeip.net>;
+ PR92881. */
+ gomp_remove_var_async (acc_dev, n, aq);
+ else
+ {
+ size_t num_mappings = 0;
+ /* If the target_mem_desc represents a single data mapping,
+ we can check that it is freed when this splay tree key's
+ refcount reaches zero. Otherwise (e.g. for a
+ 'GOMP_MAP_STRUCT' mapping with multiple members), fall
+ back to skipping the test. */
+ for (size_t l_i = 0; l_i < n->tgt->list_count; ++l_i)
+ if (n->tgt->list[l_i].key)
+ ++num_mappings;
+ bool is_tgt_unmapped = gomp_remove_var (acc_dev, n);
+ assert (is_tgt_unmapped || num_mappings > 1);
+ }
+ }
}
break;
@@ -1177,7 +1207,29 @@ goacc_exit_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum,
&& str->refcount != REFCOUNT_INFINITY)
str->refcount--;
if (str->refcount == 0)
- gomp_remove_var_async (acc_dev, str, aq);
+ {
+ if (aq)
+ /* TODO We can't do the 'is_tgt_unmapped' checking --
+ see the 'gomp_unref_tgt' comment in
+ <http://mid.mail-archive.com/878snl36eu.fsf@euler.schwinge.homeip.net>;
+ PR92881. */
+ gomp_remove_var_async (acc_dev, str, aq);
+ else
+ {
+ size_t num_mappings = 0;
+ /* If the target_mem_desc represents a single data
+ mapping, we can check that it is freed when this
+ splay tree key's refcount reaches zero.
+ Otherwise (e.g. for a 'GOMP_MAP_STRUCT' mapping
+ with multiple members), fall back to skipping
+ the test. */
+ for (size_t l_i = 0; l_i < str->tgt->list_count; ++l_i)
+ if (str->tgt->list[l_i].key)
+ ++num_mappings;
+ bool is_tgt_unmapped = gomp_remove_var (acc_dev, str);
+ assert (is_tgt_unmapped || num_mappings > 1);
+ }
+ }
}
}
i += elems;
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-refcount-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-refcount-1.c
new file mode 100644
index 000000000000..bde5890d6676
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-refcount-1.c
@@ -0,0 +1,47 @@
+/* Test dynamic unmapping of separate structure members. */
+
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+
+#include <assert.h>
+#include <stdbool.h>
+#include <openacc.h>
+
+struct s
+{
+ char a;
+ float b;
+};
+
+void test (bool use_directives)
+{
+ struct s s;
+
+#pragma acc enter data create(s.a, s.b)
+ assert (acc_is_present (&s.a, sizeof s.a));
+ assert (acc_is_present (&s.b, sizeof s.b));
+
+ if (use_directives)
+ {
+#pragma acc exit data delete(s.a)
+ }
+ else
+ acc_delete (&s.a, sizeof s.a);
+ assert (!acc_is_present (&s.a, sizeof s.a));
+ assert (acc_is_present (&s.b, sizeof s.b));
+ if (use_directives)
+ {
+#pragma acc exit data delete(s.b)
+ }
+ else
+ acc_delete (&s.b, sizeof s.b);
+ assert (!acc_is_present (&s.a, sizeof s.a));
+ assert (!acc_is_present (&s.b, sizeof s.b));
+}
+
+int main ()
+{
+ test (true);
+ test (false);
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-6.f90 b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-6.f90
index a7943d93d542..5837a4039106 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-6.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-6.f90
@@ -40,19 +40,20 @@ program dtype
if (.not. acc_is_present(var%a(5:n - 5))) stop 11
if (.not. acc_is_present(var%b(5:n - 5))) stop 12
if (.not. acc_is_present(var)) stop 13
-!$acc exit data copyout(var%a(5:n - 5), var%b(5:n - 5)) finalize
- if (acc_get_device_type() .ne. acc_device_host) then
- if (acc_is_present(var%a(5:n - 5))) stop 21
- if (acc_is_present(var%b(5:n - 5))) stop 22
- end if
print *, "CheCKpOInT1"
! { dg-output ".*CheCKpOInT1(\n|\r\n|\r)" }
- if (.not. acc_is_present(var)) stop 23
- !TODO { dg-output "STOP 23(\n|\r\n|\r)$" { target { ! openacc_host_selected } } } ! Scan for what we expect in the "XFAILed" case (without actually XFAILing).
+!$acc exit data copyout(var%a(5:n - 5), var%b(5:n - 5)) 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_get_device_type() .ne. acc_device_host) then
+ if (acc_is_present(var%a(5:n - 5))) stop 21
+ if (acc_is_present(var%b(5:n - 5))) stop 22
+ end if
+ if (.not. acc_is_present(var)) stop 23
!$acc end data
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 449a6cf6894e..445cbabb8ca4 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,16 +21,17 @@ program main
if (.not. acc_is_present(var%a)) stop 1
if (.not. acc_is_present(var)) stop 2
- !$acc exit data delete(var%a) finalize
- if (acc_is_present(var%a)) stop 3
print *, "CheCKpOInT1"
! { dg-output ".*CheCKpOInT1(\n|\r\n|\r)" }
- if (.not. acc_is_present(var)) stop 4
- !TODO { dg-output "STOP 4(\n|\r\n|\r)$" { target { ! openacc_host_selected } } } ! Scan for what we expect in the "XFAILed" case (without actually XFAILing).
+ !$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
!$acc end data
if (acc_is_present(var%a)) stop 5
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 a7e649d30415..8554534b2f27 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,16 +23,17 @@ program main
if (.not. acc_is_present(var%a)) stop 1
if (.not. acc_is_present(var)) stop 2
- !$acc exit data delete(var%a) finalize
- if (acc_is_present(var%a)) stop 3
print *, "CheCKpOInT1"
! { dg-output ".*CheCKpOInT1(\n|\r\n|\r)" }
- if (.not. acc_is_present(var)) stop 4
- !TODO { dg-output "STOP 4(\n|\r\n|\r)$" { target { ! openacc_host_selected } } } ! Scan for what we expect in the "XFAILed" case (without actually XFAILing).
+ !$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
!$acc end data
if (acc_is_present(var%a)) stop 5
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 3402fafc7e28..8e696cc70e80 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,16 +23,17 @@ program main
if (.not. acc_is_present(var%a)) stop 1
if (.not. acc_is_present(var)) stop 2
- !$acc exit data delete(var%a)
- if (acc_is_present(var%a)) stop 3
print *, "CheCKpOInT1"
! { dg-output ".*CheCKpOInT1(\n|\r\n|\r)" }
- if (.not. acc_is_present(var)) stop 4
- !TODO { dg-output "STOP 4(\n|\r\n|\r)$" { target { ! openacc_host_selected } } } ! Scan for what we expect in the "XFAILed" case (without actually XFAILing).
+ !$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
!$acc end data
if (acc_is_present(var%a)) stop 5
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 7504969d9a56..070a6f8e1493 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,16 +24,17 @@ program main
if (.not. acc_is_present(var)) stop 2
!$acc exit data detach(var%a)
- !$acc exit data delete(var%a) finalize
- if (acc_is_present(var%a)) stop 3
print *, "CheCKpOInT1"
! { dg-output ".*CheCKpOInT1(\n|\r\n|\r)" }
- if (.not. acc_is_present(var)) stop 4
- !TODO { dg-output "STOP 4(\n|\r\n|\r)$" { target { ! openacc_host_selected } } } ! Scan for what we expect in the "XFAILed" case (without actually XFAILing).
+ !$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
!$acc end data
if (acc_is_present(var%a)) stop 5
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 fedae0db054b..b22e411567ff 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,15 +23,16 @@ 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 delete(var%a)
- !TODO { dg-output "(\n|\r\n|\r)libgomp: struct not mapped for detach operation(\n|\r\n|\r)$" { target { ! openacc_host_selected } } } ! Scan for what we expect in the "XFAILed" case (without actually XFAILing).
+ !$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).
!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.26.2
next prev parent reply other threads:[~2020-06-04 18:35 UTC|newest]
Thread overview: 81+ messages / expand[flat|nested] mbox.gz Atom feed top
2018-11-10 17:11 [PATCH 0/3] OpenACC 2.6 manual deep copy support (attach/detach) Julian Brown
2018-11-10 17:11 ` [PATCH 2/3] Factor out duplicate code in gimplify_scan_omp_clauses Julian Brown
2018-12-18 14:16 ` Julian Brown
2018-12-18 14:50 ` Jakub Jelinek
2018-11-10 17:11 ` [PATCH 1/3] Host-to-device transfer coalescing & magic offset value self-documentation Julian Brown
2018-12-21 10:56 ` libgomp/target.c magic constants self-documentation Thomas Schwinge
2019-05-29 14:48 ` Thomas Schwinge
2018-11-10 17:12 ` [PATCH 3/3] OpenACC 2.6 manual deep copy support (attach/detach) Julian Brown
2018-11-11 17:04 ` Bernhard Reutner-Fischer
2018-11-30 11:41 ` [PATCH] " Julian Brown
2018-12-03 17:03 ` Julian Brown
2018-12-07 13:50 ` Jakub Jelinek
2018-12-10 19:42 ` Julian Brown
2018-12-13 10:57 ` Jakub Jelinek
2018-12-14 19:00 ` Julian Brown
2018-12-18 12:25 ` Jakub Jelinek
2018-12-22 13:37 ` Thomas Schwinge
2019-10-18 17:20 ` Thomas Schwinge
2019-11-06 18:44 ` Julian Brown
2019-11-22 23:54 ` Julian Brown
2019-11-25 10:53 ` Tobias Burnus
2019-11-26 2:54 ` Julian Brown
2019-12-17 12:16 ` Thomas Schwinge
2019-12-17 17:28 ` [WIP] OpenACC 'acc_attach*', 'acc_detach*' runtime library routines (was: [PATCH] OpenACC 2.6 manual deep copy support (attach/detach)) Thomas Schwinge
2019-12-18 6:03 ` [PATCH 00/13] OpenACC 2.6 manual deep copy support Julian Brown
2019-12-18 6:03 ` [PATCH 01/13] Use aux struct in libgomp for infrequently-used/API-specific data Julian Brown
2019-12-18 6:03 ` [PATCH 03/13] OpenACC reference count consistency checking Julian Brown
2019-12-18 6:03 ` [PATCH 02/13] OpenACC reference count overhaul Julian Brown
2020-05-19 15:42 ` Thomas Schwinge
2020-06-04 18:13 ` [OpenACC] Use 'tgt' returned from 'gomp_map_vars' (was: [PATCH 02/13] OpenACC reference count overhaul) Thomas Schwinge
2020-05-19 15:49 ` [PATCH 02/13] OpenACC reference count overhaul Thomas Schwinge
2020-05-19 15:58 ` Thomas Schwinge
2020-06-25 11:03 ` Thomas Schwinge
2020-07-03 15:29 ` Thomas Schwinge
2019-12-18 6:04 ` [PATCH 08/13] OpenACC 2.6 deep copy: middle-end parts Julian Brown
2019-12-21 21:51 ` Thomas Schwinge
2019-12-18 6:04 ` [PATCH 06/13] OpenACC 2.6 deep copy: attach/detach API routines Julian Brown
2019-12-18 6:04 ` [PATCH 09/13] OpenACC 2.6 deep copy: C and C++ front-end parts Julian Brown
2019-12-24 5:05 ` Thomas Schwinge
2019-12-26 19:04 ` Jason Merrill
2021-06-10 11:03 ` Thomas Schwinge
2019-12-18 6:04 ` [PATCH 05/13] Factor out duplicate code in gimplify_scan_omp_clauses Julian Brown
2019-12-18 6:04 ` [PATCH 04/13] Use gomp_map_val for OpenACC host-to-device address translation Julian Brown
2019-12-18 6:05 ` [PATCH 11/13] OpenACC 2.6 deep copy: C and C++ execution tests Julian Brown
2020-06-04 18:43 ` Fix 'sizeof' usage in 'libgomp.oacc-c-c++-common/deep-copy-{7, 8}.c' (was: [PATCH 11/13] OpenACC 2.6 deep copy: C and C++ execution tests) Thomas Schwinge
2023-10-31 14:00 ` Add OpenACC 'acc_map_data' variant to 'libgomp.oacc-c-c++-common/deep-copy-8.c' " Thomas Schwinge
2019-12-18 6:05 ` [PATCH 13/13] Fortran polymorphic class-type support for OpenACC Julian Brown
2019-12-18 6:05 ` [PATCH 12/13] OpenACC 2.6 deep copy: Fortran execution tests Julian Brown
2019-12-18 6:05 ` [PATCH 07/13] OpenACC 2.6 deep copy: libgomp parts Julian Brown
2019-12-21 23:37 ` Thomas Schwinge
2020-01-03 12:26 ` Julian Brown
2020-05-20 9:37 ` Thomas Schwinge
2020-06-05 16:23 ` [OpenACC 'exit data'] Simplify 'GOMP_MAP_STRUCT' handling (was: [PATCH 07/13] OpenACC 2.6 deep copy: libgomp parts) Thomas Schwinge
2020-06-05 16:36 ` [OpenACC 'exit data'] Strip 'GOMP_MAP_STRUCT' mappings " Thomas Schwinge
2020-05-20 14:52 ` [PATCH 07/13] OpenACC 2.6 deep copy: libgomp parts Thomas Schwinge
2020-05-20 19:11 ` Julian Brown
2020-06-04 18:35 ` Thomas Schwinge [this message]
2020-06-04 18:53 ` Thomas Schwinge
2020-06-05 10:39 ` Thomas Schwinge
2020-06-05 20:28 ` Julian Brown
2020-06-05 11:17 ` Thomas Schwinge
2020-06-05 20:31 ` Julian Brown
2020-06-09 10:41 ` OpenACC 'attach'/'detach' has no business affecting user-visible reference counting (was: [PATCH 07/13] OpenACC 2.6 deep copy: libgomp parts) Thomas Schwinge
2020-06-09 12:23 ` Julian Brown
2020-06-18 18:21 ` Julian Brown
2020-07-16 8:35 ` OpenACC 'attach'/'detach' has no business affecting user-visible reference counting Thomas Schwinge
2020-06-26 9:20 ` [PATCH 07/13] OpenACC 2.6 deep copy: libgomp parts Thomas Schwinge
2020-07-16 9:35 ` Thomas Schwinge
2020-07-16 21:21 ` Julian Brown
2020-07-17 9:12 ` Thomas Schwinge
2020-06-30 15:58 ` Thomas Schwinge
2019-12-18 7:20 ` [PATCH 10/13] OpenACC 2.6 deep copy: Fortran front-end parts Julian Brown
2019-12-18 23:30 ` Tobias Burnus
2019-12-20 12:25 ` [committed] Improve is-coindexed check for OpenACC/OpenMP (was: [PATCH 10/13] OpenACC 2.6 deep copy: Fortran front-end parts) Tobias Burnus
2019-12-20 13:25 ` [PATCH 10/13] OpenACC 2.6 deep copy: Fortran front-end parts Tobias Burnus
2019-12-20 10:08 ` [patch,committed] Fix testsuite-fallout of OpenACC deep-copy patch (was: [PATCH 10/13] OpenACC 2.6 deep copy: Fortran front-end parts) Tobias Burnus
2019-12-18 18:24 ` [PATCH 00/13] OpenACC 2.6 manual deep copy support Thomas Schwinge
2019-12-20 1:21 ` Julian Brown
2019-12-20 14:36 ` OpenACC regression and development pace Thomas Koenig
2020-06-04 18:07 ` [OpenACC] XFAIL behavior of over-eager 'finalize' clause (was: [PATCH 00/13] OpenACC 2.6 manual deep copy support) Thomas Schwinge
2019-12-17 16:53 ` In 'libgomp/target.c', 'struct splay_tree_key_s', use 'struct splay_tree_aux' for infrequently-used or API-specific data (was: [PATCH] OpenACC 2.6 manual deep copy support (attach/detach)) Thomas Schwinge
Reply instructions:
You may reply publicly to this message via plain-text email
using any one of the following methods:
* Save the following mbox file, import it into your mail client,
and reply-to-all from there: mbox
Avoid top-posting and favor interleaved quoting:
https://en.wikipedia.org/wiki/Posting_style#Interleaved_style
* Reply using the --to, --cc, and --in-reply-to
switches of git-send-email(1):
git send-email \
--in-reply-to=87bllysmqi.fsf@euler.schwinge.homeip.net \
--to=thomas@codesourcery.com \
--cc=gcc-patches@gcc.gnu.org \
--cc=jakub@redhat.com \
--cc=julian@codesourcery.com \
/path/to/YOUR_REPLY
https://kernel.org/pub/software/scm/git/docs/git-send-email.html
* If your mail client supports setting the In-Reply-To header
via mailto: links, try the mailto: link
Be sure your reply has a Subject: header at the top and a blank line
before the message body.
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).