* [PATCH 0/9] [OpenACC] Refcounting and manual deep copy improvements
@ 2020-06-16 22:38 Julian Brown
2020-06-16 22:38 ` [PATCH 1/9] [OpenACC] Fortran derived-type mapping fix Julian Brown
` (8 more replies)
0 siblings, 9 replies; 26+ messages in thread
From: Julian Brown @ 2020-06-16 22:38 UTC (permalink / raw)
To: gcc-patches
Cc: Thomas Schwinge, fortran, Jakub Jelinek, Tobias Burnus, Moore, Catherine
This patch series collects several improvements/bugfixes to the reference
counting and manual deep-copy implementation for OpenACC, mostly based
on review feedback from Thomas. In short, areas addressed are:
- Some unexpectedly-dead code in goacc_enter_data_internal
introduced by "Adjust dynamic reference count
semantics" is no longer dead. Flagged by Thomas in:
https://gcc.gnu.org/pipermail/gcc-patches/2020-June/547196.html
- PR95270 ("OpenACC 'enter data attach' looks up target memory object
displaced by pointer size") has been fixed for C, C++ and Fortran.
- As part of the Fortran fix above, we no longer strip GOMP_MAP_TO_PSET,
GOMP_MAP_POINTER from OpenACC "enter data" and "exit data" directives.
- We now suppress detach operations for "no-op" exit data operations
when the dynamic refcount for the associated data is zero, fixing
several newishly-added testcases.
Further discussion on individual patches. Tested (as a series) with
offloading to NVPTX. OK?
Thanks,
Julian
Julian Brown (9):
[OpenACC] Fortran derived-type mapping fix
[OpenACC] GOMP_MAP_ATTACH handling in find_group_last
[OpenACC] Adjust dynamic reference count semantics
[OpenACC] Don't pass kind array via pointer to goacc_enter_datum
[OpenACC] Fix incompatible copyout for acc_map_data (PR92843)
[OpenACC] Set bias to zero for explicit attach/detach clauses in C and
C++
[OpenACC] Do not strip GOMP_MAP_TO_PSET/GOMP_MAP_POINTER for
enter/exit data directives
[OpenACC] Fix standalone attach for Fortran assumed-shape array
pointers
[OpenACC] Don't detach for no-op exit data with zero dynamic refcount
gcc/c/c-typeck.c | 8 +
gcc/cp/semantics.c | 8 +
gcc/fortran/trans-openmp.c | 44 +++-
gcc/gimplify.c | 11 +-
gcc/testsuite/c-c++-common/goacc/mdc-1.c | 14 +-
gcc/testsuite/gfortran.dg/goacc/finalize-1.f | 4 +-
.../gfortran.dg/goacc/mapping-tests-3.f90 | 15 ++
.../gfortran.dg/goacc/mapping-tests-4.f90 | 17 ++
libgomp/libgomp.h | 8 +-
libgomp/oacc-mem.c | 243 ++++++++++++++----
libgomp/target.c | 38 +--
.../libgomp.oacc-c-c++-common/pr92843-1.c | 1 -
.../libgomp.oacc-c-c++-common/refcounting-1.c | 31 +++
.../libgomp.oacc-c-c++-common/refcounting-2.c | 31 +++
.../attach-descriptor-1.f90 | 51 ++++
.../libgomp.oacc-fortran/deep-copy-6.f90 | 6 +-
.../mdc-refcount-1-1-1.f90 | 6 +-
.../mdc-refcount-1-1-2.F90 | 2 +-
.../mdc-refcount-1-2-1.f90 | 6 +-
.../mdc-refcount-1-2-2.f90 | 6 +-
.../mdc-refcount-1-3-1.f90 | 6 +-
.../mdc-refcount-1-3-2.f90 | 5 +-
.../mdc-refcount-1-4-1.f90 | 6 +-
.../mdc-refcount-1-4-2.f90 | 5 +-
24 files changed, 430 insertions(+), 142 deletions(-)
create mode 100644 gcc/testsuite/gfortran.dg/goacc/mapping-tests-3.f90
create mode 100644 gcc/testsuite/gfortran.dg/goacc/mapping-tests-4.f90
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/refcounting-1.c
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/refcounting-2.c
create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90
--
2.23.0
^ permalink raw reply [flat|nested] 26+ messages in thread
* [PATCH 1/9] [OpenACC] Fortran derived-type mapping fix
2020-06-16 22:38 [PATCH 0/9] [OpenACC] Refcounting and manual deep copy improvements Julian Brown
@ 2020-06-16 22:38 ` Julian Brown
2020-06-16 22:38 ` [PATCH 2/9] [OpenACC] GOMP_MAP_ATTACH handling in find_group_last Julian Brown
` (7 subsequent siblings)
8 siblings, 0 replies; 26+ messages in thread
From: Julian Brown @ 2020-06-16 22:38 UTC (permalink / raw)
To: gcc-patches
Cc: Thomas Schwinge, fortran, Jakub Jelinek, Tobias Burnus, Moore, Catherine
This is a slightly-updated version of the patch sent here, with some of
Thomas's suggestions incorporated:
https://gcc.gnu.org/pipermail/gcc-patches/2020-June/547407.html
I'm still assuming this is approved, but including for completeness.
Julian
ChangeLog
gcc/fortran/
* trans-openmp.c (gfc_trans_omp_clauses): Use 'inner' not 'decl' for
derived type members which themselves have derived types.
gcc/testsuite/
* gfortran.dg/goacc/mapping-tests-3.f90: New test.
* gfortran.dg/goacc/mapping-tests-4.f90: New test.
---
gcc/fortran/trans-openmp.c | 4 ++--
.../gfortran.dg/goacc/mapping-tests-3.f90 | 15 +++++++++++++++
.../gfortran.dg/goacc/mapping-tests-4.f90 | 17 +++++++++++++++++
3 files changed, 34 insertions(+), 2 deletions(-)
create mode 100644 gcc/testsuite/gfortran.dg/goacc/mapping-tests-3.f90
create mode 100644 gcc/testsuite/gfortran.dg/goacc/mapping-tests-4.f90
diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c
index 7e2f6256c43..02c40fdc660 100644
--- a/gcc/fortran/trans-openmp.c
+++ b/gcc/fortran/trans-openmp.c
@@ -2774,9 +2774,9 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
}
else
{
- OMP_CLAUSE_DECL (node) = decl;
+ OMP_CLAUSE_DECL (node) = inner;
OMP_CLAUSE_SIZE (node)
- = TYPE_SIZE_UNIT (TREE_TYPE (decl));
+ = TYPE_SIZE_UNIT (TREE_TYPE (inner));
}
}
else if (lastcomp->next
diff --git a/gcc/testsuite/gfortran.dg/goacc/mapping-tests-3.f90 b/gcc/testsuite/gfortran.dg/goacc/mapping-tests-3.f90
new file mode 100644
index 00000000000..890ca781967
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/mapping-tests-3.f90
@@ -0,0 +1,15 @@
+! { dg-additional-options "-fdump-tree-gimple" }
+
+subroutine foo
+ type one
+ integer i, j
+ end type
+ type two
+ type(one) A, B
+ end type
+
+ type(two) x
+
+ !$acc enter data copyin(x%A)
+! { dg-final { scan-tree-dump-times "omp target oacc_enter_exit_data map\\(struct:x \\\[len: 1\\\]\\) map\\(to:x.a \\\[len: \[0-9\]+\\\]\\)" 1 "gimple" } }
+end
diff --git a/gcc/testsuite/gfortran.dg/goacc/mapping-tests-4.f90 b/gcc/testsuite/gfortran.dg/goacc/mapping-tests-4.f90
new file mode 100644
index 00000000000..17cc4841d4e
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/mapping-tests-4.f90
@@ -0,0 +1,17 @@
+subroutine foo
+ type one
+ integer i, j
+ end type
+ type two
+ type(one) A, B
+ end type
+
+ type(two) x
+
+! This is accepted at present, although it represents a probably-unintentional
+! overlapping subcopy.
+ !$acc enter data copyin(x%A, x%A%i)
+! But this raises an error.
+ !$acc enter data copyin(x%A, x%A%i, x%A%i)
+! { dg-error ".x.a.i. appears more than once in map clauses" "" { target *-*-* } .-1 }
+end
--
2.23.0
^ permalink raw reply [flat|nested] 26+ messages in thread
* [PATCH 2/9] [OpenACC] GOMP_MAP_ATTACH handling in find_group_last
2020-06-16 22:38 [PATCH 0/9] [OpenACC] Refcounting and manual deep copy improvements Julian Brown
2020-06-16 22:38 ` [PATCH 1/9] [OpenACC] Fortran derived-type mapping fix Julian Brown
@ 2020-06-16 22:38 ` Julian Brown
2020-06-30 12:42 ` Thomas Schwinge
2020-06-16 22:38 ` [PATCH 3/9] [OpenACC] Adjust dynamic reference count semantics Julian Brown
` (6 subsequent siblings)
8 siblings, 1 reply; 26+ messages in thread
From: Julian Brown @ 2020-06-16 22:38 UTC (permalink / raw)
To: gcc-patches
Cc: Thomas Schwinge, fortran, Jakub Jelinek, Tobias Burnus, Moore, Catherine
Later patches in this series assume that GOMP_MAP_ATTACH will be grouped
together with a preceding GOMP_MAP_TO_PSET or other "to" data movement
clause, except in cases where an explicit "attach" clause is used.
This patch arranges for that to be so.
OK?
Julian
ChangeLog
libgomp/
* oacc-mem.c (find_group_last): Group data-movement clauses
(GOMP_MAP_TO_PSET, GOMP_MAP_TO, etc.) together with a subsequent
GOMP_MAP_ATTACH. Allow standalone GOMP_MAP_ATTACH also.
---
libgomp/oacc-mem.c | 22 +++++++++++++++++++---
1 file changed, 19 insertions(+), 3 deletions(-)
diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c
index 936ae649dd9..be7f8d600eb 100644
--- a/libgomp/oacc-mem.c
+++ b/libgomp/oacc-mem.c
@@ -985,9 +985,15 @@ find_group_last (int pos, size_t mapnum, size_t *sizes, unsigned short *kinds)
switch (kind0)
{
case GOMP_MAP_TO_PSET:
- while (pos + 1 < mapnum && (kinds[pos + 1] & 0xff) == GOMP_MAP_POINTER)
+ if (pos + 1 < mapnum
+ && (kinds[pos + 1] & 0xff) == GOMP_MAP_ATTACH)
+ return pos + 1;
+
+ while (pos + 1 < mapnum
+ && (kinds[pos + 1] & 0xff) == GOMP_MAP_POINTER)
pos++;
- /* We expect at least one GOMP_MAP_POINTER after a GOMP_MAP_TO_PSET. */
+ /* We expect at least one GOMP_MAP_POINTER (if not a single
+ GOMP_MAP_ATTACH) after a GOMP_MAP_TO_PSET. */
assert (pos > first_pos);
break;
@@ -1002,6 +1008,9 @@ find_group_last (int pos, size_t mapnum, size_t *sizes, unsigned short *kinds)
gomp_fatal ("unexpected mapping");
break;
+ case GOMP_MAP_ATTACH:
+ return pos;
+
default:
/* GOMP_MAP_ALWAYS_POINTER can only appear directly after some other
mapping. */
@@ -1012,9 +1021,16 @@ find_group_last (int pos, size_t mapnum, size_t *sizes, unsigned short *kinds)
return pos + 1;
}
+ /* We can have a single GOMP_MAP_ATTACH mapping after a to/from
+ mapping. */
+ if (pos + 1 < mapnum
+ && (kinds[pos + 1] & 0xff) == GOMP_MAP_ATTACH)
+ return pos + 1;
+
/* 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)
+ while (pos + 1 < mapnum
+ && (kinds[pos + 1] & 0xff) == GOMP_MAP_POINTER)
pos++;
}
--
2.23.0
^ permalink raw reply [flat|nested] 26+ messages in thread
* [PATCH 3/9] [OpenACC] Adjust dynamic reference count semantics
2020-06-16 22:38 [PATCH 0/9] [OpenACC] Refcounting and manual deep copy improvements Julian Brown
2020-06-16 22:38 ` [PATCH 1/9] [OpenACC] Fortran derived-type mapping fix Julian Brown
2020-06-16 22:38 ` [PATCH 2/9] [OpenACC] GOMP_MAP_ATTACH handling in find_group_last Julian Brown
@ 2020-06-16 22:38 ` Julian Brown
2020-06-30 13:51 ` Thomas Schwinge
2020-06-16 22:38 ` [PATCH 4/9] [OpenACC] Don't pass kind array via pointer to goacc_enter_datum Julian Brown
` (5 subsequent siblings)
8 siblings, 1 reply; 26+ messages in thread
From: Julian Brown @ 2020-06-16 22:38 UTC (permalink / raw)
To: gcc-patches
Cc: Thomas Schwinge, fortran, Jakub Jelinek, Tobias Burnus, Moore, Catherine
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?
Julian
ChangeLog
libgomp/
* libgomp.h (struct splay_tree_key_s): Change virtual_refcount to
dynamic_refcount.
(struct gomp_device_descr): Remove GOMP_MAP_VARS_OPENACC_ENTER_DATA.
* oacc-mem.c (acc_map_data): Substitute virtual_refcount for
dynamic_refcount.
(goacc_enter_datum): Adjust for dynamic_refcount semantics.
(goacc_exit_datum): Re-add some error checking. Adjust for
dynamic_refcount semantics.
(goacc_enter_data_internal): Implement "present" case of dynamic
memory-map handling here. Update "non-present" case for
dynamic_refcount semantics.
(goacc_exit_data_internal): Update for dynamic_refcount semantics.
* target.c (gomp_map_vars_internal): Remove
GOMP_MAP_VARS_OPENACC_ENTER_DATA handling. Update for dynamic_refcount
handling.
(gomp_unmap_vars_internal): Remove virtual_refcount handling.
(gomp_load_image_to_device): Substitute dynamic_refcount for
virtual_refcount.
libgomp/
* testsuite/libgomp.oacc-c-c++-common/refcounting-1.c: New test.
* testsuite/libgomp.oacc-c-c++-common/refcounting-2.c: New test.
* testsuite/libgomp.oacc-fortran/deep-copy-6.f90: Remove XFAILs.
---
libgomp/libgomp.h | 8 +-
libgomp/oacc-mem.c | 155 ++++++++++++++----
libgomp/target.c | 38 +----
.../libgomp.oacc-c-c++-common/refcounting-1.c | 31 ++++
.../libgomp.oacc-c-c++-common/refcounting-2.c | 31 ++++
.../libgomp.oacc-fortran/deep-copy-6.f90 | 6 +-
6 files changed, 201 insertions(+), 68 deletions(-)
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/refcounting-1.c
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/refcounting-2.c
diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h
index ca42e0de640..7b52ce7d5c2 100644
--- a/libgomp/libgomp.h
+++ b/libgomp/libgomp.h
@@ -1016,11 +1016,8 @@ struct splay_tree_key_s {
uintptr_t tgt_offset;
/* Reference count. */
uintptr_t refcount;
- /* Reference counts beyond those that represent genuine references in the
- linked splay tree key/target memory structures, e.g. for multiple OpenACC
- "present increment" operations (via "acc enter data") referring to the same
- host-memory block. */
- uintptr_t virtual_refcount;
+ /* Dynamic reference count. */
+ uintptr_t dynamic_refcount;
struct splay_tree_aux *aux;
};
@@ -1153,7 +1150,6 @@ struct gomp_device_descr
enum gomp_map_vars_kind
{
GOMP_MAP_VARS_OPENACC,
- GOMP_MAP_VARS_OPENACC_ENTER_DATA,
GOMP_MAP_VARS_TARGET,
GOMP_MAP_VARS_DATA,
GOMP_MAP_VARS_ENTER_DATA
diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c
index be7f8d600eb..bc64bebe6c1 100644
--- a/libgomp/oacc-mem.c
+++ b/libgomp/oacc-mem.c
@@ -409,7 +409,7 @@ acc_map_data (void *h, void *d, size_t s)
splay_tree_key n = tgt->list[0].key;
assert (n);
assert (n->refcount == 1);
- assert (n->virtual_refcount == 0);
+ assert (n->dynamic_refcount == 0);
/* Special reference counting behavior. */
n->refcount = REFCOUNT_INFINITY;
@@ -456,7 +456,7 @@ acc_unmap_data (void *h)
(void *) n->host_start, (int) host_size, (void *) h);
}
/* TODO This currently doesn't catch 'REFCOUNT_INFINITY' usage different from
- 'acc_map_data'. Maybe 'virtual_refcount' can be used for disambiguating
+ 'acc_map_data'. Maybe 'dynamic_refcount' can be used for disambiguating
the different 'REFCOUNT_INFINITY' cases, or simply separate
'REFCOUNT_INFINITY' values per different usage ('REFCOUNT_ACC_MAP_DATA'
etc.)? */
@@ -545,10 +545,8 @@ goacc_enter_datum (void **hostaddrs, size_t *sizes, void *kinds, int async)
assert (n->refcount != REFCOUNT_LINK);
if (n->refcount != REFCOUNT_INFINITY)
- {
- n->refcount++;
- n->virtual_refcount++;
- }
+ n->refcount++;
+ n->dynamic_refcount++;
gomp_mutex_unlock (&acc_dev->lock);
}
@@ -562,13 +560,14 @@ goacc_enter_datum (void **hostaddrs, size_t *sizes, void *kinds, int async)
struct target_mem_desc *tgt
= gomp_map_vars_async (acc_dev, aq, mapnum, hostaddrs, NULL, sizes,
- kinds, true, GOMP_MAP_VARS_OPENACC_ENTER_DATA);
+ kinds, true, GOMP_MAP_VARS_ENTER_DATA);
assert (tgt);
assert (tgt->list_count == 1);
n = tgt->list[0].key;
assert (n);
assert (n->refcount == 1);
- assert (n->virtual_refcount == 0);
+ assert (n->dynamic_refcount == 0);
+ n->dynamic_refcount++;
d = (void *) tgt->tgt_start;
}
@@ -689,23 +688,28 @@ goacc_exit_datum (void *h, size_t s, unsigned short kind, int async)
(void *) h, (int) s, (void *) n->host_start, (int) host_size);
}
+ assert (n->refcount != REFCOUNT_LINK);
+ if (n->refcount != REFCOUNT_INFINITY
+ && n->refcount < n->dynamic_refcount)
+ {
+ gomp_mutex_unlock (&acc_dev->lock);
+ gomp_fatal ("Dynamic reference counting assert fail\n");
+ }
+
bool finalize = (kind == GOMP_MAP_DELETE
|| kind == GOMP_MAP_FORCE_FROM);
if (finalize)
{
if (n->refcount != REFCOUNT_INFINITY)
- n->refcount -= n->virtual_refcount;
- n->virtual_refcount = 0;
+ n->refcount -= n->dynamic_refcount;
+ n->dynamic_refcount = 0;
}
-
- if (n->virtual_refcount > 0)
+ else if (n->dynamic_refcount)
{
if (n->refcount != REFCOUNT_INFINITY)
n->refcount--;
- n->virtual_refcount--;
+ n->dynamic_refcount--;
}
- else if (n->refcount > 0 && n->refcount != REFCOUNT_INFINITY)
- n->refcount--;
if (n->refcount == 0)
{
@@ -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];
+ 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);
+ }
+ 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);
+ }
+
+ 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");
+ }
+ 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++;
+ }
+ }
i = group_last;
}
@@ -1148,18 +1250,15 @@ goacc_exit_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum,
if (finalize)
{
if (n->refcount != REFCOUNT_INFINITY)
- n->refcount -= n->virtual_refcount;
- n->virtual_refcount = 0;
+ n->refcount -= n->dynamic_refcount;
+ n->dynamic_refcount = 0;
}
-
- if (n->virtual_refcount > 0)
+ else if (n->dynamic_refcount)
{
if (n->refcount != REFCOUNT_INFINITY)
n->refcount--;
- n->virtual_refcount--;
+ n->dynamic_refcount--;
}
- else if (n->refcount > 0 && n->refcount != REFCOUNT_INFINITY)
- n->refcount--;
if (copyfrom
&& (kind != GOMP_MAP_FROM || n->refcount == 0))
diff --git a/libgomp/target.c b/libgomp/target.c
index 36425477dcb..3f2becdae0e 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -666,8 +666,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
struct target_mem_desc *tgt
= gomp_malloc (sizeof (*tgt) + sizeof (tgt->list[0]) * mapnum);
tgt->list_count = mapnum;
- tgt->refcount = (pragma_kind == GOMP_MAP_VARS_ENTER_DATA
- || pragma_kind == GOMP_MAP_VARS_OPENACC_ENTER_DATA) ? 0 : 1;
+ tgt->refcount = pragma_kind == GOMP_MAP_VARS_ENTER_DATA ? 0 : 1;
tgt->device_descr = devicep;
tgt->prev = NULL;
struct gomp_coalesce_buf cbuf, *cbufp = NULL;
@@ -1094,7 +1093,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
tgt->list[i].copy_from = false;
tgt->list[i].always_copy_from = false;
tgt->list[i].do_detach
- = (pragma_kind != GOMP_MAP_VARS_OPENACC_ENTER_DATA);
+ = (pragma_kind != GOMP_MAP_VARS_ENTER_DATA);
n->refcount++;
}
else
@@ -1155,7 +1154,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
tgt->list[i].offset = 0;
tgt->list[i].length = k->host_end - k->host_start;
k->refcount = 1;
- k->virtual_refcount = 0;
+ k->dynamic_refcount = 0;
tgt->refcount++;
array->left = NULL;
array->right = NULL;
@@ -1294,20 +1293,8 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
/* If the variable from "omp target enter data" map-list was already mapped,
tgt is not needed. Otherwise tgt will be freed by gomp_unmap_vars or
gomp_exit_data. */
- if ((pragma_kind == GOMP_MAP_VARS_ENTER_DATA
- || pragma_kind == GOMP_MAP_VARS_OPENACC_ENTER_DATA)
- && tgt->refcount == 0)
- {
- /* If we're about to discard a target_mem_desc with no "structural"
- references (tgt->refcount == 0), any splay keys linked in the tgt's
- list must have their virtual refcount incremented to represent that
- "lost" reference in order to implement the semantics of the OpenACC
- "present increment" operation properly. */
- if (pragma_kind == GOMP_MAP_VARS_OPENACC_ENTER_DATA)
- for (i = 0; i < tgt->list_count; i++)
- if (tgt->list[i].key)
- tgt->list[i].key->virtual_refcount++;
-
+ if (pragma_kind == GOMP_MAP_VARS_ENTER_DATA && tgt->refcount == 0)
+ {
free (tgt);
tgt = NULL;
}
@@ -1459,14 +1446,7 @@ gomp_unmap_vars_internal (struct target_mem_desc *tgt, bool do_copyfrom,
continue;
bool do_unmap = false;
- if (k->tgt == tgt
- && k->virtual_refcount > 0
- && k->refcount != REFCOUNT_INFINITY)
- {
- k->virtual_refcount--;
- k->refcount--;
- }
- else if (k->refcount > 1 && k->refcount != REFCOUNT_INFINITY)
+ if (k->refcount > 1 && k->refcount != REFCOUNT_INFINITY)
k->refcount--;
else if (k->refcount == 1)
{
@@ -1631,7 +1611,7 @@ gomp_load_image_to_device (struct gomp_device_descr *devicep, unsigned version,
k->tgt = tgt;
k->tgt_offset = target_table[i].start;
k->refcount = REFCOUNT_INFINITY;
- k->virtual_refcount = 0;
+ k->dynamic_refcount = 0;
k->aux = NULL;
array->left = NULL;
array->right = NULL;
@@ -1665,7 +1645,7 @@ gomp_load_image_to_device (struct gomp_device_descr *devicep, unsigned version,
k->tgt = tgt;
k->tgt_offset = target_var->start;
k->refcount = is_link_var ? REFCOUNT_LINK : REFCOUNT_INFINITY;
- k->virtual_refcount = 0;
+ k->dynamic_refcount = 0;
k->aux = NULL;
array->left = NULL;
array->right = NULL;
@@ -2935,7 +2915,7 @@ omp_target_associate_ptr (const void *host_ptr, const void *device_ptr,
k->tgt = tgt;
k->tgt_offset = (uintptr_t) device_ptr + device_offset;
k->refcount = REFCOUNT_INFINITY;
- k->virtual_refcount = 0;
+ k->dynamic_refcount = 0;
k->aux = NULL;
array->left = NULL;
array->right = NULL;
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/refcounting-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/refcounting-1.c
new file mode 100644
index 00000000000..4e6d06d48d5
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/refcounting-1.c
@@ -0,0 +1,31 @@
+/* Test dynamic unmapping of separate structure members. */
+
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+
+#include <assert.h>
+#include <openacc.h>
+
+struct s
+{
+ char a;
+ char b;
+};
+
+int main ()
+{
+ 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));
+
+#pragma acc exit data delete(s.a)
+#pragma acc exit data delete(s.b)
+
+ assert (!acc_is_present (&s.a, sizeof s.a));
+ assert (!acc_is_present (&s.b, sizeof s.b));
+
+ return 0;
+}
+
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/refcounting-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/refcounting-2.c
new file mode 100644
index 00000000000..5539fd8d57f
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/refcounting-2.c
@@ -0,0 +1,31 @@
+/* Test dynamic unmapping of separate structure members. */
+
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+
+#include <assert.h>
+#include <openacc.h>
+
+struct s
+{
+ char a;
+ char b;
+};
+
+int main ()
+{
+ 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));
+
+ acc_delete (&s.a, sizeof s.a);
+ 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));
+
+ 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 5837a403910..eb7d3ca160e 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-6.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-6.f90
@@ -43,12 +43,8 @@ program dtype
print *, "CheCKpOInT1"
! { dg-output ".*CheCKpOInT1(\n|\r\n|\r)" }
!$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 } } }
+ ! { dg-output ".*CheCKpOInT2(\n|\r\n|\r)" }
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
--
2.23.0
^ permalink raw reply [flat|nested] 26+ messages in thread
* [PATCH 4/9] [OpenACC] Don't pass kind array via pointer to goacc_enter_datum
2020-06-16 22:38 [PATCH 0/9] [OpenACC] Refcounting and manual deep copy improvements Julian Brown
` (2 preceding siblings ...)
2020-06-16 22:38 ` [PATCH 3/9] [OpenACC] Adjust dynamic reference count semantics Julian Brown
@ 2020-06-16 22:38 ` Julian Brown
2020-06-16 22:39 ` [PATCH 5/9] [OpenACC] Fix incompatible copyout for acc_map_data (PR92843) Julian Brown
` (4 subsequent siblings)
8 siblings, 0 replies; 26+ messages in thread
From: Julian Brown @ 2020-06-16 22:38 UTC (permalink / raw)
To: gcc-patches
Cc: Thomas Schwinge, fortran, Jakub Jelinek, Tobias Burnus, Moore, Catherine
This is a minor cleanup for goacc_enter_datum. Unchanged from previous
posting, but including for completeness:
https://gcc.gnu.org/pipermail/gcc-patches/2020-May/546331.html
OK?
Julian
ChangeLog
libgomp/
* oacc-mem.c (goacc_enter_datum): Use scalar kind argument instead of
kinds array.
(acc_create, acc_create_async, acc_copyin, acc_copyin_async): Update
calls to goacc_enter_datum.
---
libgomp/oacc-mem.c | 17 +++++++----------
1 file changed, 7 insertions(+), 10 deletions(-)
diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c
index bc64bebe6c1..05998ebc6de 100644
--- a/libgomp/oacc-mem.c
+++ b/libgomp/oacc-mem.c
@@ -501,7 +501,8 @@ acc_unmap_data (void *h)
/* Enter dynamic mapping for a single datum. Return the device pointer. */
static void *
-goacc_enter_datum (void **hostaddrs, size_t *sizes, void *kinds, int async)
+goacc_enter_datum (void **hostaddrs, size_t *sizes, unsigned short kind,
+ int async)
{
void *d;
splay_tree_key n;
@@ -560,7 +561,7 @@ goacc_enter_datum (void **hostaddrs, size_t *sizes, void *kinds, int async)
struct target_mem_desc *tgt
= gomp_map_vars_async (acc_dev, aq, mapnum, hostaddrs, NULL, sizes,
- kinds, true, GOMP_MAP_VARS_ENTER_DATA);
+ &kind, true, GOMP_MAP_VARS_ENTER_DATA);
assert (tgt);
assert (tgt->list_count == 1);
n = tgt->list[0].key;
@@ -584,15 +585,13 @@ goacc_enter_datum (void **hostaddrs, size_t *sizes, void *kinds, int async)
void *
acc_create (void *h, size_t s)
{
- unsigned short kinds[1] = { GOMP_MAP_ALLOC };
- return goacc_enter_datum (&h, &s, &kinds, acc_async_sync);
+ return goacc_enter_datum (&h, &s, GOMP_MAP_ALLOC, acc_async_sync);
}
void
acc_create_async (void *h, size_t s, int async)
{
- unsigned short kinds[1] = { GOMP_MAP_ALLOC };
- goacc_enter_datum (&h, &s, &kinds, async);
+ goacc_enter_datum (&h, &s, GOMP_MAP_ALLOC, async);
}
/* acc_present_or_create used to be what acc_create is now. */
@@ -617,15 +616,13 @@ acc_pcreate (void *h, size_t s)
void *
acc_copyin (void *h, size_t s)
{
- unsigned short kinds[1] = { GOMP_MAP_TO };
- return goacc_enter_datum (&h, &s, &kinds, acc_async_sync);
+ return goacc_enter_datum (&h, &s, GOMP_MAP_TO, acc_async_sync);
}
void
acc_copyin_async (void *h, size_t s, int async)
{
- unsigned short kinds[1] = { GOMP_MAP_TO };
- goacc_enter_datum (&h, &s, &kinds, async);
+ goacc_enter_datum (&h, &s, GOMP_MAP_TO, async);
}
/* acc_present_or_copyin used to be what acc_copyin is now. */
--
2.23.0
^ permalink raw reply [flat|nested] 26+ messages in thread
* [PATCH 5/9] [OpenACC] Fix incompatible copyout for acc_map_data (PR92843)
2020-06-16 22:38 [PATCH 0/9] [OpenACC] Refcounting and manual deep copy improvements Julian Brown
` (3 preceding siblings ...)
2020-06-16 22:38 ` [PATCH 4/9] [OpenACC] Don't pass kind array via pointer to goacc_enter_datum Julian Brown
@ 2020-06-16 22:39 ` Julian Brown
2020-06-16 22:39 ` [PATCH 6/9] [OpenACC] Set bias to zero for explicit attach/detach clauses in C and C++ Julian Brown
` (3 subsequent siblings)
8 siblings, 0 replies; 26+ messages in thread
From: Julian Brown @ 2020-06-16 22:39 UTC (permalink / raw)
To: gcc-patches
Cc: Thomas Schwinge, fortran, Jakub Jelinek, Tobias Burnus, Moore, Catherine
This is a repost of the following:
https://gcc.gnu.org/pipermail/gcc-patches/2020-May/546333.html
I made a minor edit to the ChangeLog, but the patch is otherwise
unchanged. Including for completeness.
OK?
ChangeLog
PR libgomp/92843
libgomp/
* oacc-mem.c (goacc_exit_data_internal): Don't copyout data mapped
with acc_map_data in exit data directive.
* testsuite/libgomp.oacc-c-c++-common/pr92843-1.c: Remove XFAIL.
---
libgomp/oacc-mem.c | 1 +
libgomp/testsuite/libgomp.oacc-c-c++-common/pr92843-1.c | 1 -
2 files changed, 1 insertion(+), 1 deletion(-)
diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c
index 05998ebc6de..745cb132621 100644
--- a/libgomp/oacc-mem.c
+++ b/libgomp/oacc-mem.c
@@ -1258,6 +1258,7 @@ goacc_exit_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum,
}
if (copyfrom
+ && n->refcount != REFCOUNT_INFINITY
&& (kind != GOMP_MAP_FROM || n->refcount == 0))
gomp_copy_dev2host (acc_dev, aq, (void *) cur_node.host_start,
(void *) (n->tgt->tgt_start + n->tgt_offset
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92843-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92843-1.c
index f16c46a37bf..db5b35b08d9 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92843-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92843-1.c
@@ -1,7 +1,6 @@
/* Verify that 'acc_copyout' etc. is a no-op if there's still a structured
reference count. */
-/* { dg-xfail-run-if "TODO PR92843" { *-*-* } } */
/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
#include <assert.h>
--
2.23.0
^ permalink raw reply [flat|nested] 26+ messages in thread
* [PATCH 6/9] [OpenACC] Set bias to zero for explicit attach/detach clauses in C and C++
2020-06-16 22:38 [PATCH 0/9] [OpenACC] Refcounting and manual deep copy improvements Julian Brown
` (4 preceding siblings ...)
2020-06-16 22:39 ` [PATCH 5/9] [OpenACC] Fix incompatible copyout for acc_map_data (PR92843) Julian Brown
@ 2020-06-16 22:39 ` Julian Brown
2020-06-25 11:36 ` Thomas Schwinge
2020-06-16 22:39 ` [PATCH 7/9] [OpenACC] Do not strip GOMP_MAP_TO_PSET/GOMP_MAP_POINTER for enter/exit data directives Julian Brown
` (2 subsequent siblings)
8 siblings, 1 reply; 26+ messages in thread
From: Julian Brown @ 2020-06-16 22:39 UTC (permalink / raw)
To: gcc-patches
Cc: Thomas Schwinge, fortran, Jakub Jelinek, Tobias Burnus, Moore, Catherine
This is a fix for the pointer (or array) size inadvertently being used
for the bias of attach and detach clauses (PR95270), for C and C++.
OK?
Julian
ChangeLog
PR middle-end/95270
gcc/c/
* c-typeck.c (c_finish_omp_clauses): Set OMP_CLAUSE_SIZE (bias) to zero
for standalone attach/detach clauses.
gcc/cp/
* semantics.c (finish_omp_clauses): Likewise.
gcc/testsuite/
* c-c++-common/goacc/mdc-1.c: Update expected dump output for zero
bias.
---
gcc/c/c-typeck.c | 8 ++++++++
gcc/cp/semantics.c | 8 ++++++++
gcc/testsuite/c-c++-common/goacc/mdc-1.c | 14 +++++++-------
3 files changed, 23 insertions(+), 7 deletions(-)
diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c
index 385bf3a1c7b..134f1520239 100644
--- a/gcc/c/c-typeck.c
+++ b/gcc/c/c-typeck.c
@@ -14533,6 +14533,10 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
}
if (c_oacc_check_attachments (c))
remove = true;
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH))
+ OMP_CLAUSE_SIZE (c) = size_zero_node;
break;
}
if (t == error_mark_node)
@@ -14546,6 +14550,10 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
remove = true;
break;
}
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH))
+ OMP_CLAUSE_SIZE (c) = size_zero_node;
if (TREE_CODE (t) == COMPONENT_REF
&& OMP_CLAUSE_CODE (c) != OMP_CLAUSE__CACHE_)
{
diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c
index 64587c791c6..77e6ff7fb0d 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -7334,6 +7334,10 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
}
if (cp_oacc_check_attachments (c))
remove = true;
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH))
+ OMP_CLAUSE_SIZE (c) = size_zero_node;
break;
}
if (t == error_mark_node)
@@ -7347,6 +7351,10 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
remove = true;
break;
}
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH))
+ OMP_CLAUSE_SIZE (c) = size_zero_node;
if (REFERENCE_REF_P (t)
&& TREE_CODE (TREE_OPERAND (t, 0)) == COMPONENT_REF)
{
diff --git a/gcc/testsuite/c-c++-common/goacc/mdc-1.c b/gcc/testsuite/c-c++-common/goacc/mdc-1.c
index fb5841a709d..337c1f7cc77 100644
--- a/gcc/testsuite/c-c++-common/goacc/mdc-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/mdc-1.c
@@ -45,12 +45,12 @@ t1 ()
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.to:s .len: 32.." 1 "omplower" } } */
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_data map.tofrom:.z .len: 40.. map.struct:s .len: 1.. map.alloc:s.a .len: 8.. map.tofrom:._1 .len: 40.. map.attach:s.a .bias: 0.." 1 "omplower" } } */
-/* { dg-final { scan-tree-dump-times "pragma omp target oacc_parallel map.attach:s.e .bias: 8.. map.tofrom:s .len: 32" 1 "omplower" } } */
-/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.attach:a .bias: 8.." 1 "omplower" } } */
-/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.detach:a .bias: 8.." 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "pragma omp target oacc_parallel map.attach:s.e .bias: 0.. map.tofrom:s .len: 32" 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.attach:a .bias: 0.." 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.detach:a .bias: 0.." 1 "omplower" } } */
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.to:a .len: 8.." 1 "omplower" } } */
-/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.detach:s.e .bias: 8.." 1 "omplower" } } */
-/* { dg-final { scan-tree-dump-times "pragma omp target oacc_data map.attach:s.e .bias: 8.." 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.detach:s.e .bias: 0.." 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "pragma omp target oacc_data map.attach:s.e .bias: 0.." 1 "omplower" } } */
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.release:a .len: 8.." 1 "omplower" } } */
-/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data finalize map.force_detach:a .bias: 8.." 1 "omplower" } } */
-/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data finalize map.force_detach:s.a .bias: 8.." 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data finalize map.force_detach:a .bias: 0.." 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data finalize map.force_detach:s.a .bias: 0.." 1 "omplower" } } */
--
2.23.0
^ permalink raw reply [flat|nested] 26+ messages in thread
* [PATCH 7/9] [OpenACC] Do not strip GOMP_MAP_TO_PSET/GOMP_MAP_POINTER for enter/exit data directives
2020-06-16 22:38 [PATCH 0/9] [OpenACC] Refcounting and manual deep copy improvements Julian Brown
` (5 preceding siblings ...)
2020-06-16 22:39 ` [PATCH 6/9] [OpenACC] Set bias to zero for explicit attach/detach clauses in C and C++ Julian Brown
@ 2020-06-16 22:39 ` Julian Brown
2020-07-06 16:19 ` Thomas Schwinge
2020-06-16 22:39 ` [PATCH 8/9] [OpenACC] Fix standalone attach for Fortran assumed-shape array pointers Julian Brown
2020-06-16 22:39 ` [PATCH 9/9] [OpenACC] Don't detach for no-op exit data with zero dynamic refcount Julian Brown
8 siblings, 1 reply; 26+ messages in thread
From: Julian Brown @ 2020-06-16 22:39 UTC (permalink / raw)
To: gcc-patches
Cc: Thomas Schwinge, fortran, Jakub Jelinek, Tobias Burnus, Moore, Catherine
When attaching pointers in Fortran, OpenACC 2.6 specifies that a
descriptor must be copied to the target at the same time (see next
patch). That means that stripping GOMP_MAP_TO_PSET (and lesserly,
GOMP_MAP_POINTER), which was behaviour introduced by the manual deep-copy
middle-end support patch, was probably wrong.
That arguably answers some of the questions at the end of:
https://gcc.gnu.org/pipermail/gcc-patches/2020-June/547424.html
It appears that the user can (but certainly should not!) map a synthesized
array descriptor using an "enter data" operation that can go out of
scope before that data is unmapped. It would be nice to give a warning
for an attempt to do such a thing, though I have no idea if that's
possible in practice.
gcc/
* gimplify.c (gimplify_scan_omp_clauses): Do not strip
GOMP_MAP_TO_PSET/GOMP_MAP_POINTER for OpenACC enter/exit data
directives.
gcc/testsuite/
* gfortran.dg/goacc/finalize-1.f: Update expected dump output.
---
gcc/gimplify.c | 11 ++---------
gcc/testsuite/gfortran.dg/goacc/finalize-1.f | 4 ++--
2 files changed, 4 insertions(+), 11 deletions(-)
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 9851edfc4db..aa6853f0dcc 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -8767,6 +8767,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
case OMP_TARGET_DATA:
case OMP_TARGET_ENTER_DATA:
case OMP_TARGET_EXIT_DATA:
+ case OACC_ENTER_DATA:
+ case OACC_EXIT_DATA:
case OACC_HOST_DATA:
if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
|| (OMP_CLAUSE_MAP_KIND (c)
@@ -8775,15 +8777,6 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
mapped, but not the pointer to it. */
remove = true;
break;
- case OACC_ENTER_DATA:
- case OACC_EXIT_DATA:
- if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
- || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_TO_PSET
- || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
- || (OMP_CLAUSE_MAP_KIND (c)
- == GOMP_MAP_FIRSTPRIVATE_REFERENCE))
- remove = true;
- break;
default:
break;
}
diff --git a/gcc/testsuite/gfortran.dg/goacc/finalize-1.f b/gcc/testsuite/gfortran.dg/goacc/finalize-1.f
index 1e2e3e94b8a..ca642156e9f 100644
--- a/gcc/testsuite/gfortran.dg/goacc/finalize-1.f
+++ b/gcc/testsuite/gfortran.dg/goacc/finalize-1.f
@@ -21,7 +21,7 @@
!$ACC EXIT DATA FINALIZE DELETE (del_f_p(2:5))
! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:\\*\\(c_char \\*\\) parm\\.0\\.data \\\[len: \[^\\\]\]+\\\]\\) map\\(to:del_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:\\(integer\\(kind=1\\)\\\[0:\\\] \\* restrict\\) del_f_p\\.data \\\[pointer assign, bias: \\(sizetype\\) parm\\.0\\.data - \\(sizetype\\) del_f_p\\.data\\\]\\) finalize;$" 1 "original" } }
-! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(delete:MEM\\\[\\(c_char \\*\\)\[^\\\]\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(delete:MEM\\\[\\(c_char \\*\\)\[^\\\]\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) map\\(to:del_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:del_f_p\\.data \\\[pointer assign, bias: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } }
!$ACC EXIT DATA COPYOUT (cpo_r)
! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(from:cpo_r\\);$" 1 "original" } }
@@ -33,5 +33,5 @@
!$ACC EXIT DATA COPYOUT (cpo_f_p(4:10)) FINALIZE
! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(from:\\*\\(c_char \\*\\) parm\\.1\\.data \\\[len: \[^\\\]\]+\\\]\\) map\\(to:cpo_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:\\(integer\\(kind=1\\)\\\[0:\\\] \\* restrict\\) cpo_f_p\\.data \\\[pointer assign, bias: \\(sizetype\\) parm\\.1\\.data - \\(sizetype\\) cpo_f_p\\.data\\\]\\) finalize;$" 1 "original" } }
-! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(force_from:MEM\\\[\\(c_char \\*\\)\[^\\\]\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(force_from:MEM\\\[\\(c_char \\*\\)\[^\\\]\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) map\\(to:cpo_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:cpo_f_p\\.data \\\[pointer assign, bias: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } }
END SUBROUTINE f
--
2.23.0
^ permalink raw reply [flat|nested] 26+ messages in thread
* [PATCH 8/9] [OpenACC] Fix standalone attach for Fortran assumed-shape array pointers
2020-06-16 22:38 [PATCH 0/9] [OpenACC] Refcounting and manual deep copy improvements Julian Brown
` (6 preceding siblings ...)
2020-06-16 22:39 ` [PATCH 7/9] [OpenACC] Do not strip GOMP_MAP_TO_PSET/GOMP_MAP_POINTER for enter/exit data directives Julian Brown
@ 2020-06-16 22:39 ` Julian Brown
2020-07-14 11:43 ` Thomas Schwinge
2020-06-16 22:39 ` [PATCH 9/9] [OpenACC] Don't detach for no-op exit data with zero dynamic refcount Julian Brown
8 siblings, 1 reply; 26+ messages in thread
From: Julian Brown @ 2020-06-16 22:39 UTC (permalink / raw)
To: gcc-patches
Cc: Thomas Schwinge, fortran, Jakub Jelinek, Tobias Burnus, Moore, Catherine
As mentioned in the blurb for the previous patch, an "attach" operation
for a Fortran pointer with an array descriptor must copy that array
descriptor to the target. This patch arranges for that to be so.
OK?
Julian
ChangeLog
gcc/fortran/
* trans-openmp.c (gfc_trans_omp_clauses): Copy array descriptor to
target for attach clauses when appropriate.
libgomp/
* testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90: New test.
---
gcc/fortran/trans-openmp.c | 40 ++++++++++++++-
.../attach-descriptor-1.f90 | 51 +++++++++++++++++++
2 files changed, 89 insertions(+), 2 deletions(-)
create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90
diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c
index 02c40fdc660..909a86795e0 100644
--- a/gcc/fortran/trans-openmp.c
+++ b/gcc/fortran/trans-openmp.c
@@ -2573,8 +2573,44 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
}
}
if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (decl))
- && n->u.map_op != OMP_MAP_ATTACH
- && n->u.map_op != OMP_MAP_DETACH)
+ && (n->u.map_op == OMP_MAP_ATTACH
+ || n->u.map_op == OMP_MAP_DETACH))
+ {
+ tree type = TREE_TYPE (decl);
+ tree data = gfc_conv_descriptor_data_get (decl);
+ if (present)
+ data = gfc_build_cond_assign_expr (block, present,
+ data,
+ null_pointer_node);
+ tree ptr
+ = fold_convert (build_pointer_type (char_type_node),
+ data);
+ ptr = build_fold_indirect_ref (ptr);
+ /* Standalone attach clauses used with arrays with
+ descriptors must copy the descriptor to the target,
+ else they won't have anything to perform the
+ attachment onto (see OpenACC 2.6, "2.6.3. Data
+ Structures with Pointers"). */
+ OMP_CLAUSE_DECL (node) = ptr;
+ node2 = build_omp_clause (input_location, OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (node2, GOMP_MAP_TO_PSET);
+ OMP_CLAUSE_DECL (node2) = decl;
+ OMP_CLAUSE_SIZE (node2) = TYPE_SIZE_UNIT (type);
+ node3 = build_omp_clause (input_location, OMP_CLAUSE_MAP);
+ if (n->u.map_op == OMP_MAP_ATTACH)
+ {
+ OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_ATTACH);
+ n->u.map_op = OMP_MAP_ALLOC;
+ }
+ else /* OMP_MAP_DETACH. */
+ {
+ OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_DETACH);
+ n->u.map_op = OMP_MAP_RELEASE;
+ }
+ OMP_CLAUSE_DECL (node3) = data;
+ OMP_CLAUSE_SIZE (node3) = size_int (0);
+ }
+ else if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (decl)))
{
tree type = TREE_TYPE (decl);
tree ptr = gfc_conv_descriptor_data_get (decl);
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90
new file mode 100644
index 00000000000..2dd1a6fa5b6
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90
@@ -0,0 +1,51 @@
+program att
+ use openacc
+ implicit none
+ type t
+ integer :: arr1(10)
+ integer, allocatable :: arr2(:)
+ end type t
+ integer :: i
+ type(t) :: myvar
+ integer, target :: tarr(10)
+ integer, pointer :: myptr(:)
+
+ allocate(myvar%arr2(10))
+
+ do i=1,10
+ myvar%arr1(i) = 0
+ myvar%arr2(i) = 0
+ tarr(i) = 0
+ end do
+
+ call acc_copyin(myvar)
+ call acc_copyin(myvar%arr2)
+ call acc_copyin(tarr)
+
+ myptr => tarr
+
+ !$acc enter data attach(myvar%arr2, myptr)
+
+ ! FIXME: This warning is emitted on the wrong line number.
+ ! { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } 36 }
+ !$acc serial present(myvar%arr2)
+ do i=1,10
+ myvar%arr1(i) = i
+ myvar%arr2(i) = i
+ end do
+ myptr(3) = 99
+ !$acc end serial
+
+ !$acc exit data detach(myvar%arr2, myptr)
+
+ call acc_copyout(myvar%arr2)
+ call acc_copyout(myvar)
+ call acc_copyout(tarr)
+
+ do i=1,10
+ if (myvar%arr1(i) .ne. i) stop 1
+ if (myvar%arr2(i) .ne. i) stop 2
+ end do
+ if (tarr(3) .ne. 99) stop 3
+
+end program att
--
2.23.0
^ permalink raw reply [flat|nested] 26+ messages in thread
* [PATCH 9/9] [OpenACC] Don't detach for no-op exit data with zero dynamic refcount
2020-06-16 22:38 [PATCH 0/9] [OpenACC] Refcounting and manual deep copy improvements Julian Brown
` (7 preceding siblings ...)
2020-06-16 22:39 ` [PATCH 8/9] [OpenACC] Fix standalone attach for Fortran assumed-shape array pointers Julian Brown
@ 2020-06-16 22:39 ` Julian Brown
2020-07-24 14:18 ` Thomas Schwinge
8 siblings, 1 reply; 26+ messages in thread
From: Julian Brown @ 2020-06-16 22:39 UTC (permalink / raw)
To: gcc-patches
Cc: Thomas Schwinge, fortran, Jakub Jelinek, Tobias Burnus, Moore, Catherine
This patch fixes a set of XFAILs in some recently-added patches by
skipping a detach operation on "no-op" exit data operations for blocks
with zero dynamic refcount. This takes advantage of the ordering of
detach clauses with respect to associated data-movement clauses: i.e.,
they are grouped together adjacently.
OK?
Julian
ChangeLog
libgomp/
* oacc-mem.c (find_group_last): Handle detach operations.
(goacc_exit_data_internal): Detect detachments that are part of copyout
operations, and suppress them if dynamic refcount is zero.
* testsuite/libgomp.oacc-fortran/mdc-refcount-1-1-1.f90: Remove XFAILs.
* testsuite/libgomp.oacc-fortran/mdc-refcount-1-1-2.F90: Fix typo.
* testsuite/libgomp.oacc-fortran/mdc-refcount-1-2-1.f90: Remove XFAILs.
* 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-3-2.f90: Likewise.
* testsuite/libgomp.oacc-fortran/mdc-refcount-1-4-1.f90: Likewise.
* testsuite/libgomp.oacc-fortran/mdc-refcount-1-4-2.f90: Likewise.
---
libgomp/oacc-mem.c | 54 ++++++++++++++++---
.../mdc-refcount-1-1-1.f90 | 6 +--
.../mdc-refcount-1-1-2.F90 | 2 +-
.../mdc-refcount-1-2-1.f90 | 6 +--
.../mdc-refcount-1-2-2.f90 | 6 +--
.../mdc-refcount-1-3-1.f90 | 6 +--
.../mdc-refcount-1-3-2.f90 | 5 +-
.../mdc-refcount-1-4-1.f90 | 6 +--
.../mdc-refcount-1-4-2.f90 | 5 +-
9 files changed, 55 insertions(+), 41 deletions(-)
diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c
index 745cb132621..f852652c048 100644
--- a/libgomp/oacc-mem.c
+++ b/libgomp/oacc-mem.c
@@ -987,7 +987,9 @@ find_group_last (int pos, size_t mapnum, size_t *sizes, unsigned short *kinds)
{
case GOMP_MAP_TO_PSET:
if (pos + 1 < mapnum
- && (kinds[pos + 1] & 0xff) == GOMP_MAP_ATTACH)
+ && ((kinds[pos + 1] & 0xff) == GOMP_MAP_ATTACH
+ || (kinds[pos + 1] & 0xff) == GOMP_MAP_DETACH
+ || (kinds[pos + 1] & 0xff) == GOMP_MAP_FORCE_DETACH))
return pos + 1;
while (pos + 1 < mapnum
@@ -1010,6 +1012,8 @@ find_group_last (int pos, size_t mapnum, size_t *sizes, unsigned short *kinds)
break;
case GOMP_MAP_ATTACH:
+ case GOMP_MAP_DETACH:
+ case GOMP_MAP_FORCE_DETACH:
return pos;
default:
@@ -1025,7 +1029,9 @@ find_group_last (int pos, size_t mapnum, size_t *sizes, unsigned short *kinds)
/* We can have a single GOMP_MAP_ATTACH mapping after a to/from
mapping. */
if (pos + 1 < mapnum
- && (kinds[pos + 1] & 0xff) == GOMP_MAP_ATTACH)
+ && ((kinds[pos + 1] & 0xff) == GOMP_MAP_ATTACH
+ || (kinds[pos + 1] & 0xff) == GOMP_MAP_DETACH
+ || (kinds[pos + 1] & 0xff) == GOMP_MAP_FORCE_DETACH))
return pos + 1;
/* We can have zero or more GOMP_MAP_POINTER mappings after a to/from
@@ -1168,15 +1174,43 @@ goacc_exit_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum,
{
gomp_mutex_lock (&acc_dev->lock);
- /* Handle "detach" before copyback/deletion of mapped data. */
- for (size_t i = 0; i < mapnum; ++i)
+ /* Handle "detach" before copyback/deletion of mapped data. If this isn't a
+ standalone "detach" clause, take care to skip the "detach" operation if
+ the dynamic refcount of the data to be detached is zero. */
+ for (size_t grp = 0; grp < mapnum; grp++)
{
- unsigned char kind = kinds[i] & 0xff;
+ size_t i = grp, group_last = find_group_last (grp, mapnum, sizes, kinds);
+ unsigned char kind = kinds[grp] & 0xff;
bool finalize = false;
+
switch (kind)
{
+ case GOMP_MAP_TO_PSET:
+ case GOMP_MAP_TOFROM:
+ case GOMP_MAP_FROM:
+ case GOMP_MAP_FORCE_FROM:
+ case GOMP_MAP_RELEASE:
+ case GOMP_MAP_DELETE:
+ {
+ if (i + 1 >= mapnum)
+ break;
+ kind = kinds[i + 1] & 0xff;
+ if (kind != GOMP_MAP_FORCE_DETACH && kind != GOMP_MAP_DETACH)
+ break;
+ splay_tree_key n = lookup_host (acc_dev, hostaddrs[i], sizes[i]);
+ if (n == NULL)
+ {
+ gomp_mutex_unlock (&acc_dev->lock);
+ gomp_fatal ("target data not mapped for detach operation");
+ }
+ i++;
+ if (n->dynamic_refcount == 0)
+ break;
+ }
+ /* Fallthrough. */
+
case GOMP_MAP_FORCE_DETACH:
- finalize = true;
+ finalize = (kind == GOMP_MAP_FORCE_DETACH);
/* Fallthrough. */
case GOMP_MAP_DETACH:
@@ -1197,9 +1231,15 @@ goacc_exit_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum,
gomp_detach_pointer (acc_dev, aq, n, hostaddr, finalize, NULL);
}
break;
+ case GOMP_MAP_STRUCT:
+ case GOMP_MAP_POINTER:
+ /* Ignore. */
+ break;
default:
- ;
+ gomp_fatal (">>>> goacc_exit_data_internal UNHANDLED kind 0x%.2x",
+ kind);
}
+ grp = group_last;
}
for (size_t i = 0; i < mapnum; ++i)
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..7171affb9f0 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
@@ -24,12 +24,8 @@ program main
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 } } }
+ ! { dg-output ".*CheCKpOInT2(\n|\r\n|\r)" }
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..2aa46189e9a 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
@@ -6,4 +6,4 @@
#include "mdc-refcount-1-1-1.f90"
! { dg-output ".*CheCKpOInT1(\n|\r\n|\r)" }
-! { dg-output ".CheCKpOInT2(\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..9a10aa5a781 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
@@ -26,12 +26,8 @@ program main
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 } } }
+ ! { dg-output ".*CheCKpOInT2(\n|\r\n|\r)" }
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..f506adf8e91 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
@@ -26,12 +26,8 @@ program main
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 } } }
+ ! { dg-output ".*CheCKpOInT2(\n|\r\n|\r)" }
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..450d95d3686 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
@@ -27,12 +27,8 @@ program main
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 } } }
+ ! { dg-output ".*CheCKpOInT2(\n|\r\n|\r)" }
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-2.f90 b/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-3-2.f90
index 3c4bbda7f66..35efad4138a 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-3-2.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-3-2.f90
@@ -27,11 +27,8 @@ program main
print *, "CheCKpOInT1"
! { dg-output ".*CheCKpOInT1(\n|\r\n|\r)" }
!$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 } } }
+ ! { dg-output ".*CheCKpOInT2(\n|\r\n|\r)" }
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..816562fc055 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
@@ -26,12 +26,8 @@ program main
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).
- !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 } } }
+ ! { dg-output ".*CheCKpOInT2(\n|\r\n|\r)" }
!$acc exit data delete(var%a)
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-2.f90 b/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-4-2.f90
index 476cd5c1bee..b98bfd74924 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-4-2.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-4-2.f90
@@ -27,11 +27,8 @@ program main
print *, "CheCKpOInT1"
! { dg-output ".*CheCKpOInT1(\n|\r\n|\r)" }
!$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 } } }
+ ! { dg-output ".*CheCKpOInT2(\n|\r\n|\r)" }
if (acc_is_present(var%a)) stop 3
if (.not. acc_is_present(var)) stop 4
--
2.23.0
^ permalink raw reply [flat|nested] 26+ messages in thread
* Re: [PATCH 6/9] [OpenACC] Set bias to zero for explicit attach/detach clauses in C and C++
2020-06-16 22:39 ` [PATCH 6/9] [OpenACC] Set bias to zero for explicit attach/detach clauses in C and C++ Julian Brown
@ 2020-06-25 11:36 ` Thomas Schwinge
2020-07-09 21:06 ` Thomas Schwinge
0 siblings, 1 reply; 26+ messages in thread
From: Thomas Schwinge @ 2020-06-25 11:36 UTC (permalink / raw)
To: Julian Brown
Cc: gcc-patches, Jakub Jelinek, Tobias Burnus, Catherine_Moore, fortran
[-- Attachment #1: Type: text/plain, Size: 4978 bytes --]
Hi Julian!
On 2020-06-16T15:39:42-0700, Julian Brown <julian@codesourcery.com> wrote:
> This is a fix for the pointer (or array) size inadvertently being used
> for the bias of attach and detach clauses (PR95270)
Thanks for looking into that one, which had caused my some gray hair.
> for C and C++.
That means, there is no such problem for Fortran? (I haven't run into
one, just curious.)
> OK?
In principle, yes, for master and releases/gcc-10 branches, but please
incorporate the following items:
> PR middle-end/95270
>
> gcc/c/
> * c-typeck.c (c_finish_omp_clauses): Set OMP_CLAUSE_SIZE (bias) to zero
> for standalone attach/detach clauses.
>
> gcc/cp/
> * semantics.c (finish_omp_clauses): Likewise.
>
> gcc/testsuite/
> * c-c++-common/goacc/mdc-1.c: Update expected dump output for zero
> bias.
> ---
> gcc/c/c-typeck.c | 8 ++++++++
> gcc/cp/semantics.c | 8 ++++++++
> gcc/testsuite/c-c++-common/goacc/mdc-1.c | 14 +++++++-------
> 3 files changed, 23 insertions(+), 7 deletions(-)
> --- a/gcc/c/c-typeck.c
> +++ b/gcc/c/c-typeck.c
> @@ -14533,6 +14533,10 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
> }
> if (c_oacc_check_attachments (c))
> remove = true;
> + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
> + && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
> + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH))
> + OMP_CLAUSE_SIZE (c) = size_zero_node;
> break;
> }
> if (t == error_mark_node)
> @@ -14546,6 +14550,10 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
> remove = true;
> break;
> }
> + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
> + && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
> + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH))
> + OMP_CLAUSE_SIZE (c) = size_zero_node;
> if (TREE_CODE (t) == COMPONENT_REF
> && OMP_CLAUSE_CODE (c) != OMP_CLAUSE__CACHE_)
> {
I cannot comment if these two code paths are good places (and the only
ones) that need to set 'OMP_CLAUSE_SIZE', so I'll trust you've found the
best/all places.
Does that override an 'OMP_CLAUSE_SIZE' that was set earlier, or
initialize it? Maybe the latter, given my comment in
<https://gcc.gnu.org/PR95270>: "make sure to skip/invalidate the
'gcc/gimplify.c:gimplify_scan_omp_clauses' handling"?
Plase add some commentary here in the code, instead of just in the
ChangeLog, something like: "initialize here, so that gimplify doesn't
wrongly do so later" (if that's what it is, and in proper language, of
course).
> --- a/gcc/cp/semantics.c
> +++ b/gcc/cp/semantics.c
> @@ -7334,6 +7334,10 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
> }
> if (cp_oacc_check_attachments (c))
> remove = true;
> + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
> + && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
> + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH))
> + OMP_CLAUSE_SIZE (c) = size_zero_node;
> break;
> }
> if (t == error_mark_node)
> @@ -7347,6 +7351,10 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
> remove = true;
> break;
> }
> + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
> + && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
> + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH))
> + OMP_CLAUSE_SIZE (c) = size_zero_node;
> if (REFERENCE_REF_P (t)
> && TREE_CODE (TREE_OPERAND (t, 0)) == COMPONENT_REF)
> {
Likewise.
> --- a/gcc/testsuite/c-c++-common/goacc/mdc-1.c
> +++ b/gcc/testsuite/c-c++-common/goacc/mdc-1.c
Obvious.
In <https://gcc.gnu.org/PR95270> I also requested size vs. bias be
documented in 'include/gomp-constants.h:enum gomp_map_kind'.
Generally, I'm still somewhat confused by the 'bias' usage in libgomp.
Is it really only used for the *initial* attach, but then (correctly so?)
ignored for any later ones? Please add some commentary next to the
respective libgomp code.
Please also include an execution test case, like I had included with
<https://gcc.gnu.org/PR95270>, for example the two files I'm attaching.
Ah actually, since the directive variant now no longer fails, please
merge these into one file, with 'test(bool directive)', and two
'test(false)', 'test(true)' calls from 'main'.
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: pr95270_-d.c --]
[-- Type: text/x-csrc, Size: 42 bytes --]
#define DIRECTIVE
#include "pr95270_-r.c"
[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #3: pr95270_-r.c --]
[-- Type: text/x-csrc, Size: 449 bytes --]
// <https://gcc.gnu.org/PR95270>
#include <assert.h>
#include <openacc.h>
int main()
{
int data;
int *data_p_dev = (int *) acc_create(&data, sizeof data);
int *data_p = &data;
acc_copyin(&data_p, sizeof data_p);
#ifdef DIRECTIVE
# pragma acc enter data attach(data_p)
#else
{
void **ptr = (void **) &data_p;
acc_attach(ptr);
}
#endif
acc_update_self(&data_p, sizeof data_p);
assert (data_p == data_p_dev);
return 0;
}
^ permalink raw reply [flat|nested] 26+ messages in thread
* Re: [PATCH 2/9] [OpenACC] GOMP_MAP_ATTACH handling in find_group_last
2020-06-16 22:38 ` [PATCH 2/9] [OpenACC] GOMP_MAP_ATTACH handling in find_group_last Julian Brown
@ 2020-06-30 12:42 ` Thomas Schwinge
0 siblings, 0 replies; 26+ messages in thread
From: Thomas Schwinge @ 2020-06-30 12:42 UTC (permalink / raw)
To: Julian Brown
Cc: gcc-patches, Jakub Jelinek, Tobias Burnus, Moore, Catherine, fortran
Hi Julian!
On 2020-06-16T15:38:32-0700, Julian Brown <julian@codesourcery.com> wrote:
> Later patches in this series assume that GOMP_MAP_ATTACH will be grouped
> together with a preceding GOMP_MAP_TO_PSET or other "to" data movement
> clause, except in cases where an explicit "attach" clause is used.
> This patch arranges for that to be so.
>
> OK?
OK for master branch and releases/gcc-10 branch. However, still a few
questions, which can be addressed separately:
> --- a/libgomp/oacc-mem.c
> +++ b/libgomp/oacc-mem.c
> @@ -985,9 +985,15 @@ find_group_last (int pos, size_t mapnum, size_t *sizes, unsigned short *kinds)
> switch (kind0)
> {
> case GOMP_MAP_TO_PSET:
> - while (pos + 1 < mapnum && (kinds[pos + 1] & 0xff) == GOMP_MAP_POINTER)
> + if (pos + 1 < mapnum
> + && (kinds[pos + 1] & 0xff) == GOMP_MAP_ATTACH)
> + return pos + 1;
> +
> + while (pos + 1 < mapnum
> + && (kinds[pos + 1] & 0xff) == GOMP_MAP_POINTER)
> pos++;
> - /* We expect at least one GOMP_MAP_POINTER after a GOMP_MAP_TO_PSET. */
> + /* We expect at least one GOMP_MAP_POINTER (if not a single
> + GOMP_MAP_ATTACH) after a GOMP_MAP_TO_PSET. */
> assert (pos > first_pos);
> break;
So a 'GOMP_MAP_TO_PSET' can now be followed by either a single
'GOMP_MAP_ATTACH', or by potentially several 'GOMP_MAP_POINTER's, but not
both. If the former ('GOMP_MAP_ATTACH'), then any additional following
'GOMP_MAP_POINTER's are not anymore handled together with the
'GOMP_MAP_TO_PSET' -- which defeats the description in
'include/gomp-constants.h', which explicitly details how
'GOMP_MAP_TO_PSET' is used to specially handle 'GOMP_MAP_POINTER's
following it. So, please update the 'enum gomp_map_kind' definition to
describe what's (now) actually going on.
(Maybe that'll then make obsolete the source code comments you're adding
here? ..., if also updating the 'GOMP_MAP_ATTACH' description to detail
how it may follow 'GOMP_MAP_TO' etc. I think such rationale --
describing the valid combinations -- is better put there instead of into
'find_group_last'.)
In the compiler, are we making sure that after 'GOMP_MAP_TO_PSET' we're
not trying to emit both a 'GOMP_MAP_ATTACH' as well as
'GOMP_MAP_POINTER's?
> @@ -1002,6 +1008,9 @@ find_group_last (int pos, size_t mapnum, size_t *sizes, unsigned short *kinds)
> gomp_fatal ("unexpected mapping");
> break;
>
> + case GOMP_MAP_ATTACH:
> + return pos;
> +
> default:
> /* GOMP_MAP_ALWAYS_POINTER can only appear directly after some other
> mapping. */
> @@ -1012,9 +1021,16 @@ find_group_last (int pos, size_t mapnum, size_t *sizes, unsigned short *kinds)
> return pos + 1;
> }
>
> + /* We can have a single GOMP_MAP_ATTACH mapping after a to/from
> + mapping. */
> + if (pos + 1 < mapnum
> + && (kinds[pos + 1] & 0xff) == GOMP_MAP_ATTACH)
> + return pos + 1;
> +
> /* 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)
> + while (pos + 1 < mapnum
> + && (kinds[pos + 1] & 0xff) == GOMP_MAP_POINTER)
> pos++;
> }
Similar (regarding documenting in 'enum gomp_map_kind' etc.).
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
^ permalink raw reply [flat|nested] 26+ messages in thread
* Re: [PATCH 3/9] [OpenACC] Adjust dynamic reference count semantics
2020-06-16 22:38 ` [PATCH 3/9] [OpenACC] Adjust dynamic reference count semantics Julian Brown
@ 2020-06-30 13:51 ` Thomas Schwinge
2020-07-03 15:41 ` Thomas Schwinge
0 siblings, 1 reply; 26+ messages in thread
From: Thomas Schwinge @ 2020-06-30 13:51 UTC (permalink / raw)
To: Julian Brown
Cc: gcc-patches, Jakub Jelinek, Tobias Burnus, Moore, Catherine, fortran
[-- Attachment #1: Type: text/plain, Size: 7241 bytes --]
Hi Julian!
On 2020-06-16T15:38:33-0700, Julian Brown <julian@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
<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,
<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 <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 <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
[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #2: 0001-into-Adjust-dynamic-reference-count-semantics.patch --]
[-- Type: text/x-diff, Size: 8424 bytes --]
From 31f7b23a8ec4107898d612f2c758f39faa0f0691 Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <thomas@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
[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #3: struct-3-1-1.c --]
[-- Type: text/x-csrc, Size: 1301 bytes --]
/* 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;
}
^ permalink raw reply [flat|nested] 26+ messages in thread
* Re: [PATCH 3/9] [OpenACC] Adjust dynamic reference count semantics
2020-06-30 13:51 ` Thomas Schwinge
@ 2020-07-03 15:41 ` Thomas Schwinge
2020-07-10 12:08 ` Julian Brown
0 siblings, 1 reply; 26+ messages in thread
From: Thomas Schwinge @ 2020-07-03 15:41 UTC (permalink / raw)
To: Julian Brown
Cc: gcc-patches, Jakub Jelinek, Tobias Burnus, Catherine_Moore, fortran
[-- Attachment #1: Type: text/plain, Size: 2545 bytes --]
Hi Julian!
On 2020-06-30T15:51:14+0200, I wrote:
> On 2020-06-16T15:38:33-0700, Julian Brown <julian@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?)
Ah, you said "Tested (as a series)", so that's probably why I saw this
intermediate step but you didn't.
> 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".
Given my recent "[OpenACC] Revert always-copyfrom behavior for
'GOMP_MAP_FORCE_FROM' in 'libgomp/oacc-mem.c:goacc_exit_data_internal'",
<http://mid.mail-archive.com/87wo3ky5vn.fsf@euler.schwinge.homeip.net>,
please also include the attached "into 'Adjust dynamic reference count
semantics': un-XFAIL 'libgomp.oacc-c-c++-common/pr92843-1.c'".
> Your patch regresses the attached
> 'libgomp.oacc-c-c++-common/struct-3-1-1.c'
That was confusing: that's a new test case, not yet in tree.
> 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
[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #2: 0001-into-Adjust-dynamic-reference-count-semantics-un-XFA.patch --]
[-- Type: text/x-diff, Size: 1922 bytes --]
From 705898afc94c94545a2dd7ed9f451615c067385f Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <thomas@codesourcery.com>
Date: Fri, 3 Jul 2020 16:58:34 +0200
Subject: [PATCH] into 'Adjust dynamic reference count semantics': un-XFAIL
'libgomp.oacc-c-c++-common/pr92843-1.c'
---
libgomp/testsuite/libgomp.oacc-c-c++-common/pr92843-1.c | 9 ---------
1 file changed, 9 deletions(-)
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92843-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92843-1.c
index 78fe1402ad46..db5b35b08d9f 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92843-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92843-1.c
@@ -4,7 +4,6 @@
/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
#include <assert.h>
-#include <stdio.h>
#include <stdlib.h>
#include <openacc.h>
@@ -135,15 +134,7 @@ test_acc_data ()
assert (acc_is_present (h, sizeof h));
assign_array (h, N, c1);
- fprintf (stderr, "CheCKpOInT1\n");
- // { dg-output ".*CheCKpOInT1(\n|\r\n|\r)" }
acc_copyout_finalize (h, sizeof h);
- //TODO goacc_exit_datum: 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.
- fprintf (stderr, "CheCKpOInT2\n");
- // { dg-output ".CheCKpOInT2(\n|\r\n|\r)" { target { openacc_host_selected } } }
assert (acc_is_present (h, sizeof h));
verify_array (h, N, c1);
--
2.27.0
^ permalink raw reply [flat|nested] 26+ messages in thread
* Re: [PATCH 7/9] [OpenACC] Do not strip GOMP_MAP_TO_PSET/GOMP_MAP_POINTER for enter/exit data directives
2020-06-16 22:39 ` [PATCH 7/9] [OpenACC] Do not strip GOMP_MAP_TO_PSET/GOMP_MAP_POINTER for enter/exit data directives Julian Brown
@ 2020-07-06 16:19 ` Thomas Schwinge
0 siblings, 0 replies; 26+ messages in thread
From: Thomas Schwinge @ 2020-07-06 16:19 UTC (permalink / raw)
To: Julian Brown
Cc: gcc-patches, Jakub Jelinek, Tobias Burnus, Catherine_Moore, fortran
[-- Attachment #1: Type: text/plain, Size: 6705 bytes --]
Hi Julian!
On 2020-06-16T15:39:43-0700, Julian Brown <julian@codesourcery.com> wrote:
> When attaching pointers in Fortran, OpenACC 2.6 specifies that a
> descriptor must be copied to the target at the same time (see next
> patch). That means that stripping GOMP_MAP_TO_PSET (and lesserly,
> GOMP_MAP_POINTER), which was behaviour introduced by the manual deep-copy
> middle-end support patch, was probably wrong.
>
> That arguably answers some of the questions at the end of:
>
> https://gcc.gnu.org/pipermail/gcc-patches/2020-June/547424.html
ACK.
> It appears that the user can (but certainly should not!) map a synthesized
> array descriptor using an "enter data" operation that can go out of
> scope before that data is unmapped. It would be nice to give a warning
> for an attempt to do such a thing, though I have no idea if that's
> possible in practice.
That's a rather complex scenario. ;-)
If I'm understanding this right, what we need to show is that an object
is created as a persistent, visible device copy, with state initialized
by 'enter data', and then any 'GOMP_MAP_TO_PSET' etc. that come with each
OpenACC 'parallel' etc. are no-ops (because the object is present
already).
My attached (new) 'libgomp.oacc-fortran/dynamic-pointer-1.f90' would seem
to be a conceptually simple test case for this, using a Fortran
'pointer'. (I hope I got my Fortran right, please verify.) This test
case doesn't work in current master and releases/gcc-10 branches (because
we don't create the persistent, visible device copy), and is "enabled" by
your patch posted here. I'm intentionally not saying "regression fixed"
or something like that, because it also doesn't work before all the
"OpenACC 2.6 deep copy: middle-end parts" etc. changes... (Maybe because
of wrong handling of 'GOMP_MAP_TO_PSET' back then, too? Just mentioning
that for completeness; I don't think we need to investigate that now.)
Please include some such rationale in the commit log, or "even" as source
code comments, as makes sense. This code surely is complicated/complex
to grasp.
> gcc/
> * gimplify.c (gimplify_scan_omp_clauses): Do not strip
> GOMP_MAP_TO_PSET/GOMP_MAP_POINTER for OpenACC enter/exit data
> directives.
Should reference PR92929 here.
Please include my attached (new)
'libgomp.oacc-fortran/dynamic-pointer-1.f90' (assuming that one makes
sense to you), and then this is OK for master and releases/gcc-10
branches.
We then (later) still need to resolve other items discussed in PR92929
"OpenACC/OpenMP 'target' 'exit data'/'update' optimizations".
Grüße
Thomas
> gcc/gimplify.c | 11 ++---------
> gcc/testsuite/gfortran.dg/goacc/finalize-1.f | 4 ++--
> 2 files changed, 4 insertions(+), 11 deletions(-)
>
> diff --git a/gcc/gimplify.c b/gcc/gimplify.c
> index 9851edfc4db..aa6853f0dcc 100644
> --- a/gcc/gimplify.c
> +++ b/gcc/gimplify.c
> @@ -8767,6 +8767,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
> case OMP_TARGET_DATA:
> case OMP_TARGET_ENTER_DATA:
> case OMP_TARGET_EXIT_DATA:
> + case OACC_ENTER_DATA:
> + case OACC_EXIT_DATA:
> case OACC_HOST_DATA:
> if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
> || (OMP_CLAUSE_MAP_KIND (c)
> @@ -8775,15 +8777,6 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
> mapped, but not the pointer to it. */
> remove = true;
> break;
> - case OACC_ENTER_DATA:
> - case OACC_EXIT_DATA:
> - if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
> - || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_TO_PSET
> - || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
> - || (OMP_CLAUSE_MAP_KIND (c)
> - == GOMP_MAP_FIRSTPRIVATE_REFERENCE))
> - remove = true;
> - break;
> default:
> break;
> }
> diff --git a/gcc/testsuite/gfortran.dg/goacc/finalize-1.f b/gcc/testsuite/gfortran.dg/goacc/finalize-1.f
> index 1e2e3e94b8a..ca642156e9f 100644
> --- a/gcc/testsuite/gfortran.dg/goacc/finalize-1.f
> +++ b/gcc/testsuite/gfortran.dg/goacc/finalize-1.f
> @@ -21,7 +21,7 @@
>
> !$ACC EXIT DATA FINALIZE DELETE (del_f_p(2:5))
> ! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:\\*\\(c_char \\*\\) parm\\.0\\.data \\\[len: \[^\\\]\]+\\\]\\) map\\(to:del_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:\\(integer\\(kind=1\\)\\\[0:\\\] \\* restrict\\) del_f_p\\.data \\\[pointer assign, bias: \\(sizetype\\) parm\\.0\\.data - \\(sizetype\\) del_f_p\\.data\\\]\\) finalize;$" 1 "original" } }
> -! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(delete:MEM\\\[\\(c_char \\*\\)\[^\\\]\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } }
> +! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(delete:MEM\\\[\\(c_char \\*\\)\[^\\\]\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) map\\(to:del_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:del_f_p\\.data \\\[pointer assign, bias: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } }
>
> !$ACC EXIT DATA COPYOUT (cpo_r)
> ! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(from:cpo_r\\);$" 1 "original" } }
> @@ -33,5 +33,5 @@
>
> !$ACC EXIT DATA COPYOUT (cpo_f_p(4:10)) FINALIZE
> ! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(from:\\*\\(c_char \\*\\) parm\\.1\\.data \\\[len: \[^\\\]\]+\\\]\\) map\\(to:cpo_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:\\(integer\\(kind=1\\)\\\[0:\\\] \\* restrict\\) cpo_f_p\\.data \\\[pointer assign, bias: \\(sizetype\\) parm\\.1\\.data - \\(sizetype\\) cpo_f_p\\.data\\\]\\) finalize;$" 1 "original" } }
> -! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(force_from:MEM\\\[\\(c_char \\*\\)\[^\\\]\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } }
> +! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(force_from:MEM\\\[\\(c_char \\*\\)\[^\\\]\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) map\\(to:cpo_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:cpo_f_p\\.data \\\[pointer assign, bias: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } }
> END SUBROUTINE f
-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter
[-- Attachment #2: dynamic-pointer-1.f90 --]
[-- Type: text/plain, Size: 2325 bytes --]
! Verify that a 'enter data'ed 'pointer' object creates a persistent, visible device copy
! { dg-do run }
! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } }
module m
implicit none
contains
subroutine verify_a (a_ref, a)
implicit none
integer, dimension (:, :, :), allocatable :: a_ref
integer, dimension (:, :, :), pointer :: a
!$acc routine seq
if (any (lbound (a) /= lbound (a_ref))) stop 101
if (any (ubound (a) /= ubound (a_ref))) stop 102
if (size (a) /= size (a_ref)) stop 103
end subroutine verify_a
end module m
program main
use m
use openacc
implicit none
integer, parameter :: n = 30
integer, dimension (:, :, :), allocatable, target :: a1, a2
integer, dimension (:, :, :), pointer :: p
allocate (a1(1:n, 0:n-1, 10:n/2))
!$acc enter data create(a1)
allocate (a2(3:n/3, 10:n, n-10:n+10))
!$acc enter data create(a2)
p => a1
call verify_a(a1, p)
! 'p' object isn't present on the device.
!$acc parallel ! Implicit 'copy(p)'; creates 'p' object...
call verify_a(a1, p)
!$acc end parallel ! ..., and deletes it again.
p => a2
call verify_a(a2, p)
! 'p' object isn't present on the device.
!$acc parallel ! Implicit 'copy(p)'; creates 'p' object...
call verify_a(a2, p)
!$acc end parallel ! ..., and deletes it again.
p => a1
!$acc enter data create(p)
! 'p' object is now present on the device (visible device copy).
!TODO PR96080 if (.not. acc_is_present (p)) stop 1
!$acc parallel
! On the device, got created as 'p => a1'.
call verify_a(a1, p)
!$acc end parallel
call verify_a(a1, p)
!$acc parallel
p => a2
! On the device, 'p => a2' is now set.
call verify_a(a2, p)
!$acc end parallel
! On the host, 'p => a1' persists.
call verify_a(a1, p)
!$acc parallel
! On the device, 'p => a2' persists.
call verify_a(a2, p)
!$acc end parallel
! On the host, 'p => a1' still persists.
call verify_a(a1, p)
p => a2
!$acc parallel
p => a1
! On the device, 'p => a1' is now set.
call verify_a(a1, p)
!$acc end parallel
! On the host, 'p => a2' persists.
call verify_a(a2, p)
!$acc parallel
! On the device, 'p => a1' persists.
call verify_a(a1, p)
!$acc end parallel
! On the host, 'p => a2' still persists.
call verify_a(a2, p)
end program main
^ permalink raw reply [flat|nested] 26+ messages in thread
* Re: [PATCH 6/9] [OpenACC] Set bias to zero for explicit attach/detach clauses in C and C++
2020-06-25 11:36 ` Thomas Schwinge
@ 2020-07-09 21:06 ` Thomas Schwinge
2020-07-09 21:32 ` Julian Brown
0 siblings, 1 reply; 26+ messages in thread
From: Thomas Schwinge @ 2020-07-09 21:06 UTC (permalink / raw)
To: Julian Brown
Cc: gcc-patches, Jakub Jelinek, Tobias Burnus, Catherine_Moore, fortran
Hi Julian!
On 2020-06-25T13:36:15+0200, I wrote:
> On 2020-06-16T15:39:42-0700, Julian Brown <julian@codesourcery.com> wrote:
>> This is a fix for the pointer (or array) size inadvertently being used
>> for the bias of attach and detach clauses (PR95270)
>
> Thanks for looking into that one, which had caused my some gray hair.
>
>> for C and C++.
>
> That means, there is no such problem for Fortran? (I haven't run into
> one, just curious.)
Looking into something else, I've now found the very same (?) problem for
Fortran, too. :-| For the following simple testcase, I again do see
non-zero 'bias: 64' for 'enter data attach(data_p)':
program main
use openacc
implicit none
!TODO Per PR96080, data types chosen so that we can create a "pointer object 'data_p'" on the device.
integer, dimension(:), target :: data(1)
integer, dimension(:), pointer :: data_p
!TODO Per PR96080, not using OpenACC/Fortran runtime library routines.
!$acc enter data create(data)
data_p => data
!$acc enter data copyin(data_p)
!$acc enter data attach(data_p)
end program main
..., and the 'attach' fails with 'libgomp: pointer target not mapped for
attach'. It doesn't fail when I force 'bias = 0' in
'gomp_attach_pointer'.
I've tried a bit, but it seems a bit difficult in Fortran to verify (with
'associated(data_p, data)' etc.) what we've actually 'attach'ed: per
PR96080, a 'call acc_update_self(data_p)' may not be doing the expected
thing, and a '!$acc update self(data_p)' per
'libgomp/oacc-parallel.c:GOACC_update' will update the actual data, but
is no-op for 'GOMP_MAP_TO_PSET', 'GOMP_MAP_POINTER'. I've stopped
experimenting with that any further.
So it seems Fortran front end changes will also be required in addition
to the C, C++ front end changes you've come up with. (For avoidance of
doubt: OK to do separately, if you'd like to. Please also reference GCC
PR95270 for these, and include the testcase from above, or something
better.)
Grüße
Thomas
> In principle, yes, for master and releases/gcc-10 branches, but please
> incorporate the following items:
>
>> PR middle-end/95270
>>
>> gcc/c/
>> * c-typeck.c (c_finish_omp_clauses): Set OMP_CLAUSE_SIZE (bias) to zero
>> for standalone attach/detach clauses.
>>
>> gcc/cp/
>> * semantics.c (finish_omp_clauses): Likewise.
>>
>> gcc/testsuite/
>> * c-c++-common/goacc/mdc-1.c: Update expected dump output for zero
>> bias.
>> ---
>> gcc/c/c-typeck.c | 8 ++++++++
>> gcc/cp/semantics.c | 8 ++++++++
>> gcc/testsuite/c-c++-common/goacc/mdc-1.c | 14 +++++++-------
>> 3 files changed, 23 insertions(+), 7 deletions(-)
>
>> --- a/gcc/c/c-typeck.c
>> +++ b/gcc/c/c-typeck.c
>> @@ -14533,6 +14533,10 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
>> }
>> if (c_oacc_check_attachments (c))
>> remove = true;
>> + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
>> + && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
>> + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH))
>> + OMP_CLAUSE_SIZE (c) = size_zero_node;
>> break;
>> }
>> if (t == error_mark_node)
>> @@ -14546,6 +14550,10 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
>> remove = true;
>> break;
>> }
>> + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
>> + && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
>> + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH))
>> + OMP_CLAUSE_SIZE (c) = size_zero_node;
>> if (TREE_CODE (t) == COMPONENT_REF
>> && OMP_CLAUSE_CODE (c) != OMP_CLAUSE__CACHE_)
>> {
>
> I cannot comment if these two code paths are good places (and the only
> ones) that need to set 'OMP_CLAUSE_SIZE', so I'll trust you've found the
> best/all places.
>
> Does that override an 'OMP_CLAUSE_SIZE' that was set earlier, or
> initialize it? Maybe the latter, given my comment in
> <https://gcc.gnu.org/PR95270>: "make sure to skip/invalidate the
> 'gcc/gimplify.c:gimplify_scan_omp_clauses' handling"?
>
> Plase add some commentary here in the code, instead of just in the
> ChangeLog, something like: "initialize here, so that gimplify doesn't
> wrongly do so later" (if that's what it is, and in proper language, of
> course).
>
>> --- a/gcc/cp/semantics.c
>> +++ b/gcc/cp/semantics.c
>> @@ -7334,6 +7334,10 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
>> }
>> if (cp_oacc_check_attachments (c))
>> remove = true;
>> + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
>> + && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
>> + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH))
>> + OMP_CLAUSE_SIZE (c) = size_zero_node;
>> break;
>> }
>> if (t == error_mark_node)
>> @@ -7347,6 +7351,10 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
>> remove = true;
>> break;
>> }
>> + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
>> + && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
>> + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH))
>> + OMP_CLAUSE_SIZE (c) = size_zero_node;
>> if (REFERENCE_REF_P (t)
>> && TREE_CODE (TREE_OPERAND (t, 0)) == COMPONENT_REF)
>> {
>
> Likewise.
>
>> --- a/gcc/testsuite/c-c++-common/goacc/mdc-1.c
>> +++ b/gcc/testsuite/c-c++-common/goacc/mdc-1.c
>
> Obvious.
>
> In <https://gcc.gnu.org/PR95270> I also requested size vs. bias be
> documented in 'include/gomp-constants.h:enum gomp_map_kind'.
>
> Generally, I'm still somewhat confused by the 'bias' usage in libgomp.
> Is it really only used for the *initial* attach, but then (correctly so?)
> ignored for any later ones? Please add some commentary next to the
> respective libgomp code.
>
> Please also include an execution test case, like I had included with
> <https://gcc.gnu.org/PR95270>, for example the two files I'm attaching.
> Ah actually, since the directive variant now no longer fails, please
> merge these into one file, with 'test(bool directive)', and two
> 'test(false)', 'test(true)' calls from 'main'.
>
>
> Grüße
> Thomas
> [ pr95270_-d.c: text/x-csrc ]
> #define DIRECTIVE
> #include "pr95270_-r.c"
> [ pr95270_-r.c: text/x-csrc ]
> // <https://gcc.gnu.org/PR95270>
>
> #include <assert.h>
> #include <openacc.h>
>
> int main()
> {
> int data;
> int *data_p_dev = (int *) acc_create(&data, sizeof data);
> int *data_p = &data;
> acc_copyin(&data_p, sizeof data_p);
>
> #ifdef DIRECTIVE
> # pragma acc enter data attach(data_p)
> #else
> {
> void **ptr = (void **) &data_p;
> acc_attach(ptr);
> }
> #endif
>
> acc_update_self(&data_p, sizeof data_p);
> assert (data_p == data_p_dev);
>
> return 0;
> }
-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter
^ permalink raw reply [flat|nested] 26+ messages in thread
* Re: [PATCH 6/9] [OpenACC] Set bias to zero for explicit attach/detach clauses in C and C++
2020-07-09 21:06 ` Thomas Schwinge
@ 2020-07-09 21:32 ` Julian Brown
0 siblings, 0 replies; 26+ messages in thread
From: Julian Brown @ 2020-07-09 21:32 UTC (permalink / raw)
To: Thomas Schwinge
Cc: gcc-patches, Jakub Jelinek, Tobias Burnus, Catherine_Moore, fortran
On Thu, 9 Jul 2020 23:06:29 +0200
Thomas Schwinge <thomas@codesourcery.com> wrote:
> Hi Julian!
>
> On 2020-06-25T13:36:15+0200, I wrote:
> > On 2020-06-16T15:39:42-0700, Julian Brown <julian@codesourcery.com>
> > wrote:
> >> This is a fix for the pointer (or array) size inadvertently being
> >> used for the bias of attach and detach clauses (PR95270)
> >
> > Thanks for looking into that one, which had caused my some gray
> > hair.
> >> for C and C++.
> >
> > That means, there is no such problem for Fortran? (I haven't run
> > into one, just curious.)
>
> Looking into something else, I've now found the very same (?) problem
> for Fortran, too. :-| For the following simple testcase, I again do
> see non-zero 'bias: 64' for 'enter data attach(data_p)':
>
> program main
> use openacc
> implicit none
> !TODO Per PR96080, data types chosen so that we can create a
> "pointer object 'data_p'" on the device. integer, dimension(:),
> target :: data(1) integer, dimension(:), pointer :: data_p
>
> !TODO Per PR96080, not using OpenACC/Fortran runtime library
> routines.
> !$acc enter data create(data)
> data_p => data
> !$acc enter data copyin(data_p)
>
> !$acc enter data attach(data_p)
> end program main
>
> ..., and the 'attach' fails with 'libgomp: pointer target not mapped
> for attach'. It doesn't fail when I force 'bias = 0' in
> 'gomp_attach_pointer'.
>
> I've tried a bit, but it seems a bit difficult in Fortran to verify
> (with 'associated(data_p, data)' etc.) what we've actually
> 'attach'ed: per PR96080, a 'call acc_update_self(data_p)' may not be
> doing the expected thing, and a '!$acc update self(data_p)' per
> 'libgomp/oacc-parallel.c:GOACC_update' will update the actual data,
> but is no-op for 'GOMP_MAP_TO_PSET', 'GOMP_MAP_POINTER'. I've stopped
> experimenting with that any further.
>
> So it seems Fortran front end changes will also be required in
> addition to the C, C++ front end changes you've come up with. (For
> avoidance of doubt: OK to do separately, if you'd like to. Please
> also reference GCC PR95270 for these, and include the testcase from
> above, or something better.)
Do the 7th & 8th patches in this series help? They were "supposed to"
be the Fortran equivalent of these C/C++ changes, though I found
additional problems too.
Thanks,
Julian
^ permalink raw reply [flat|nested] 26+ messages in thread
* Re: [PATCH 3/9] [OpenACC] Adjust dynamic reference count semantics
2020-07-03 15:41 ` Thomas Schwinge
@ 2020-07-10 12:08 ` Julian Brown
0 siblings, 0 replies; 26+ messages in thread
From: Julian Brown @ 2020-07-10 12:08 UTC (permalink / raw)
To: Thomas Schwinge
Cc: gcc-patches, Jakub Jelinek, Tobias Burnus, Catherine_Moore, fortran
On Fri, 3 Jul 2020 17:41:12 +0200
Thomas Schwinge <thomas@codesourcery.com> wrote:
> Hi Julian!
>
> On 2020-06-30T15:51:14+0200, I wrote:
> > On 2020-06-16T15:38:33-0700, Julian Brown <julian@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?)
>
> Ah, you said "Tested (as a series)", so that's probably why I saw this
> intermediate step but you didn't.
>
> > 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".
>
> Given my recent "[OpenACC] Revert always-copyfrom behavior for
> 'GOMP_MAP_FORCE_FROM' in
> 'libgomp/oacc-mem.c:goacc_exit_data_internal'",
> <http://mid.mail-archive.com/87wo3ky5vn.fsf@euler.schwinge.homeip.net>,
> please also include the attached "into 'Adjust dynamic reference
> count semantics': un-XFAIL 'libgomp.oacc-c-c++-common/pr92843-1.c'".
>
>
> > Your patch regresses the attached
> > 'libgomp.oacc-c-c++-common/struct-3-1-1.c'
>
> That was confusing: that's a new test case, not yet in tree.
I've posted a new version of the patch here that (hopefully!) addresses
all review comments:
https://gcc.gnu.org/pipermail/gcc-patches/2020-July/549774.html
Thanks,
Julian
^ permalink raw reply [flat|nested] 26+ messages in thread
* Re: [PATCH 8/9] [OpenACC] Fix standalone attach for Fortran assumed-shape array pointers
2020-06-16 22:39 ` [PATCH 8/9] [OpenACC] Fix standalone attach for Fortran assumed-shape array pointers Julian Brown
@ 2020-07-14 11:43 ` Thomas Schwinge
2020-07-15 10:28 ` Thomas Schwinge
0 siblings, 1 reply; 26+ messages in thread
From: Thomas Schwinge @ 2020-07-14 11:43 UTC (permalink / raw)
To: Julian Brown, Tobias Burnus
Cc: gcc-patches, Jakub Jelinek, Catherine_Moore, fortran
Hi Julian, Tobias!
On 2020-06-16T15:39:44-0700, Julian Brown <julian@codesourcery.com> wrote:
> As mentioned in the blurb for the previous patch, an "attach" operation
> for a Fortran pointer with an array descriptor must copy that array
> descriptor to the target.
Heh, I see -- I don't think I had read the OpenACC standard in that way,
but I think I agree your interpretation is fine.
This does not create some sort of memory leak -- everything implicitly
allocated there will eventually be deallocated again, right?
> This patch arranges for that to be so.
In response to the new OpenACC/Fortran testcase that I'd submtited in
<http://mid.mail-archive.com/87wo3co0tm.fsf@euler.schwinge.homeip.net>,
you (Julian) correctly supposed in
<http://mid.mail-archive.com/20200709223246.23a4d0e0@squid.athome>, that
this patch indeed does resolve that testcase, too. That wasn't obvious
to me. So, similar to
'libgomp/testsuite/libgomp.oacc-c-c++-common/pr95270-{1.2}.c', please
include my new OpenACC/Fortran testcase (if that makes sense to you), and
reference PR95270 in the commit log.
> OK?
Basically yes (for master and releases/gcc-10 branches), but please
consider the following:
> --- a/gcc/fortran/trans-openmp.c
> +++ b/gcc/fortran/trans-openmp.c
> @@ -2573,8 +2573,44 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
> }
> }
> if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (decl))
> - && n->u.map_op != OMP_MAP_ATTACH
> - && n->u.map_op != OMP_MAP_DETACH)
> + && (n->u.map_op == OMP_MAP_ATTACH
> + || n->u.map_op == OMP_MAP_DETACH))
> + {
> + tree type = TREE_TYPE (decl);
> + tree data = gfc_conv_descriptor_data_get (decl);
> + if (present)
> + data = gfc_build_cond_assign_expr (block, present,
> + data,
> + null_pointer_node);
> + tree ptr
> + = fold_convert (build_pointer_type (char_type_node),
> + data);
> + ptr = build_fold_indirect_ref (ptr);
> + /* Standalone attach clauses used with arrays with
> + descriptors must copy the descriptor to the target,
> + else they won't have anything to perform the
> + attachment onto (see OpenACC 2.6, "2.6.3. Data
> + Structures with Pointers"). */
> + OMP_CLAUSE_DECL (node) = ptr;
> + node2 = build_omp_clause (input_location, OMP_CLAUSE_MAP);
> + OMP_CLAUSE_SET_MAP_KIND (node2, GOMP_MAP_TO_PSET);
> + OMP_CLAUSE_DECL (node2) = decl;
> + OMP_CLAUSE_SIZE (node2) = TYPE_SIZE_UNIT (type);
> + node3 = build_omp_clause (input_location, OMP_CLAUSE_MAP);
> + if (n->u.map_op == OMP_MAP_ATTACH)
> + {
> + OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_ATTACH);
> + n->u.map_op = OMP_MAP_ALLOC;
> + }
> + else /* OMP_MAP_DETACH. */
> + {
> + OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_DETACH);
> + n->u.map_op = OMP_MAP_RELEASE;
> + }
> + OMP_CLAUSE_DECL (node3) = data;
> + OMP_CLAUSE_SIZE (node3) = size_int (0);
> + }
So this ("case A") duplicates most of the code from...
> + else if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (decl)))
> {
> [...]
... this existing case here ("case B"). It's not clear to me if these
two cases really still need to be handled separately, and a little bit
differently (regarding 'if (present)' handling, for example), or if they
could/should (?) be merged? Tobias, do you have an opinion?
Do we have sufficient testsuite coverage? (For example,
'attach'/'detach' with 'present == false', if that makes sense, or any
other thing that case A is doing differently from case B?) Shouldn't
this get '-fdump-tree-original' and/or '-fdump-tree-gimple' testcases,
similar to 'gfortran.dg/goacc/finalize-1.f', so that we verify/document
what we generate here?
> --- /dev/null
> +++ b/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90
> @@ -0,0 +1,51 @@
> +program att
Please add 'dg-do run', and...
> + use openacc
> + implicit none
> + type t
> + integer :: arr1(10)
> + integer, allocatable :: arr2(:)
> + end type t
> + integer :: i
> + type(t) :: myvar
> + integer, target :: tarr(10)
> + integer, pointer :: myptr(:)
> +
> + allocate(myvar%arr2(10))
> +
> + do i=1,10
> + myvar%arr1(i) = 0
> + myvar%arr2(i) = 0
> + tarr(i) = 0
> + end do
> +
> + call acc_copyin(myvar)
> + call acc_copyin(myvar%arr2)
> + call acc_copyin(tarr)
> +
> + myptr => tarr
> +
> + !$acc enter data attach(myvar%arr2, myptr)
> +
> + ! FIXME: This warning is emitted on the wrong line number.
> + ! { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } 36 }
... don't forget to adjust "36" here. ;-)
> + !$acc serial present(myvar%arr2)
> + do i=1,10
> + myvar%arr1(i) = i
> + myvar%arr2(i) = i
> + end do
> + myptr(3) = 99
> + !$acc end serial
> +
> + !$acc exit data detach(myvar%arr2, myptr)
> +
> + call acc_copyout(myvar%arr2)
> + call acc_copyout(myvar)
> + call acc_copyout(tarr)
> +
> + do i=1,10
> + if (myvar%arr1(i) .ne. i) stop 1
> + if (myvar%arr2(i) .ne. i) stop 2
> + end do
> + if (tarr(3) .ne. 99) stop 3
> +
> +end program att
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
^ permalink raw reply [flat|nested] 26+ messages in thread
* Re: [PATCH 8/9] [OpenACC] Fix standalone attach for Fortran assumed-shape array pointers
2020-07-14 11:43 ` Thomas Schwinge
@ 2020-07-15 10:28 ` Thomas Schwinge
2020-07-17 11:16 ` Thomas Schwinge
0 siblings, 1 reply; 26+ messages in thread
From: Thomas Schwinge @ 2020-07-15 10:28 UTC (permalink / raw)
To: Julian Brown, Tobias Burnus
Cc: gcc-patches, Jakub Jelinek, Catherine_Moore, fortran
Hi Julian, Tobias!
On 2020-07-14T13:43:37+0200, I wrote:
> On 2020-06-16T15:39:44-0700, Julian Brown <julian@codesourcery.com> wrote:
>> As mentioned in the blurb for the previous patch, an "attach" operation
>> for a Fortran pointer with an array descriptor must copy that array
>> descriptor to the target.
>
> Heh, I see -- I don't think I had read the OpenACC standard in that way,
> but I think I agree your interpretation is fine.
>
> This does not create some sort of memory leak -- everything implicitly
> allocated there will eventually be deallocated again, right?
>
>> This patch arranges for that to be so.
>
> In response to the new OpenACC/Fortran testcase that I'd submtited in
> <http://mid.mail-archive.com/87wo3co0tm.fsf@euler.schwinge.homeip.net>,
> you (Julian) correctly supposed in
> <http://mid.mail-archive.com/20200709223246.23a4d0e0@squid.athome>, that
> this patch indeed does resolve that testcase, too. That wasn't obvious
> to me. So, similar to
> 'libgomp/testsuite/libgomp.oacc-c-c++-common/pr95270-{1.2}.c', please
> include my new OpenACC/Fortran testcase (if that makes sense to you), and
> reference PR95270 in the commit log.
My new OpenACC/Fortran testcase got again broken ('libgomp: pointer
target not mapped for attach') by Tobias' commit
102502e32ea4e8a75d6b252ba319d09d735d9aa7 "[OpenMP, Fortran] Add
structure/derived-type element mapping",
<http://mid.mail-archive.com/c5b43e02-d1d5-e7cf-c11c-6daf1e8f33c5@codesourcery.com>.
Similar ('libgomp: attempt to attach null pointer') for your new
'libgomp.oacc-fortran/attach-descriptor-1.f90'.
(Whether or not 'attach'ing 'NULL' should actually be allowed, is a
separate topic for discussion.)
So this patch here will (obviously) need to be adapted to what Tobias
changed. (Plus my more general questions quoted above and below.)
Grüße
Thomas
>> OK?
>
> Basically yes (for master and releases/gcc-10 branches), but please
> consider the following:
>
>> --- a/gcc/fortran/trans-openmp.c
>> +++ b/gcc/fortran/trans-openmp.c
>> @@ -2573,8 +2573,44 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
>> }
>> }
>> if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (decl))
>> - && n->u.map_op != OMP_MAP_ATTACH
>> - && n->u.map_op != OMP_MAP_DETACH)
>> + && (n->u.map_op == OMP_MAP_ATTACH
>> + || n->u.map_op == OMP_MAP_DETACH))
>> + {
>> + tree type = TREE_TYPE (decl);
>> + tree data = gfc_conv_descriptor_data_get (decl);
>> + if (present)
>> + data = gfc_build_cond_assign_expr (block, present,
>> + data,
>> + null_pointer_node);
>> + tree ptr
>> + = fold_convert (build_pointer_type (char_type_node),
>> + data);
>> + ptr = build_fold_indirect_ref (ptr);
>> + /* Standalone attach clauses used with arrays with
>> + descriptors must copy the descriptor to the target,
>> + else they won't have anything to perform the
>> + attachment onto (see OpenACC 2.6, "2.6.3. Data
>> + Structures with Pointers"). */
>> + OMP_CLAUSE_DECL (node) = ptr;
>> + node2 = build_omp_clause (input_location, OMP_CLAUSE_MAP);
>> + OMP_CLAUSE_SET_MAP_KIND (node2, GOMP_MAP_TO_PSET);
>> + OMP_CLAUSE_DECL (node2) = decl;
>> + OMP_CLAUSE_SIZE (node2) = TYPE_SIZE_UNIT (type);
>> + node3 = build_omp_clause (input_location, OMP_CLAUSE_MAP);
>> + if (n->u.map_op == OMP_MAP_ATTACH)
>> + {
>> + OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_ATTACH);
>> + n->u.map_op = OMP_MAP_ALLOC;
>> + }
>> + else /* OMP_MAP_DETACH. */
>> + {
>> + OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_DETACH);
>> + n->u.map_op = OMP_MAP_RELEASE;
>> + }
>> + OMP_CLAUSE_DECL (node3) = data;
>> + OMP_CLAUSE_SIZE (node3) = size_int (0);
>> + }
>
> So this ("case A") duplicates most of the code from...
>
>> + else if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (decl)))
>> {
>> [...]
>
> ... this existing case here ("case B"). It's not clear to me if these
> two cases really still need to be handled separately, and a little bit
> differently (regarding 'if (present)' handling, for example), or if they
> could/should (?) be merged? Tobias, do you have an opinion?
>
> Do we have sufficient testsuite coverage? (For example,
> 'attach'/'detach' with 'present == false', if that makes sense, or any
> other thing that case A is doing differently from case B?) Shouldn't
> this get '-fdump-tree-original' and/or '-fdump-tree-gimple' testcases,
> similar to 'gfortran.dg/goacc/finalize-1.f', so that we verify/document
> what we generate here?
>
>
>> --- /dev/null
>> +++ b/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90
>> @@ -0,0 +1,51 @@
>> +program att
>
> Please add 'dg-do run', and...
>
>> + use openacc
>> + implicit none
>> + type t
>> + integer :: arr1(10)
>> + integer, allocatable :: arr2(:)
>> + end type t
>> + integer :: i
>> + type(t) :: myvar
>> + integer, target :: tarr(10)
>> + integer, pointer :: myptr(:)
>> +
>> + allocate(myvar%arr2(10))
>> +
>> + do i=1,10
>> + myvar%arr1(i) = 0
>> + myvar%arr2(i) = 0
>> + tarr(i) = 0
>> + end do
>> +
>> + call acc_copyin(myvar)
>> + call acc_copyin(myvar%arr2)
>> + call acc_copyin(tarr)
>> +
>> + myptr => tarr
>> +
>> + !$acc enter data attach(myvar%arr2, myptr)
>> +
>> + ! FIXME: This warning is emitted on the wrong line number.
>> + ! { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } 36 }
>
> ... don't forget to adjust "36" here. ;-)
>
>> + !$acc serial present(myvar%arr2)
>> + do i=1,10
>> + myvar%arr1(i) = i
>> + myvar%arr2(i) = i
>> + end do
>> + myptr(3) = 99
>> + !$acc end serial
>> +
>> + !$acc exit data detach(myvar%arr2, myptr)
>> +
>> + call acc_copyout(myvar%arr2)
>> + call acc_copyout(myvar)
>> + call acc_copyout(tarr)
>> +
>> + do i=1,10
>> + if (myvar%arr1(i) .ne. i) stop 1
>> + if (myvar%arr2(i) .ne. i) stop 2
>> + end do
>> + if (tarr(3) .ne. 99) stop 3
>> +
>> +end program att
>
>
> 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
^ permalink raw reply [flat|nested] 26+ messages in thread
* Re: [PATCH 8/9] [OpenACC] Fix standalone attach for Fortran assumed-shape array pointers
2020-07-15 10:28 ` Thomas Schwinge
@ 2020-07-17 11:16 ` Thomas Schwinge
2020-07-27 14:33 ` Julian Brown
0 siblings, 1 reply; 26+ messages in thread
From: Thomas Schwinge @ 2020-07-17 11:16 UTC (permalink / raw)
To: Julian Brown, Tobias Burnus
Cc: gcc-patches, Jakub Jelinek, Catherine_Moore, fortran
[-- Attachment #1: Type: text/plain, Size: 10031 bytes --]
Hi Julian, Tobias!
On 2020-07-15T12:28:42+0200, Thomas Schwinge <thomas@codesourcery.com> wrote:
> On 2020-07-14T13:43:37+0200, I wrote:
>> On 2020-06-16T15:39:44-0700, Julian Brown <julian@codesourcery.com> wrote:
>>> As mentioned in the blurb for the previous patch, an "attach" operation
>>> for a Fortran pointer with an array descriptor must copy that array
>>> descriptor to the target.
>>
>> Heh, I see -- I don't think I had read the OpenACC standard in that way,
>> but I think I agree your interpretation is fine.
>>
>> This does not create some sort of memory leak -- everything implicitly
>> allocated there will eventually be deallocated again, right?
Unanswered -- but I may now have found this problem, and also found "the
reverse problem" ('finalize'); see below.
>>> This patch arranges for that to be so.
>>
>> In response to the new OpenACC/Fortran testcase that I'd submtited in
>> <http://mid.mail-archive.com/87wo3co0tm.fsf@euler.schwinge.homeip.net>,
>> you (Julian) correctly supposed in
>> <http://mid.mail-archive.com/20200709223246.23a4d0e0@squid.athome>, that
>> this patch indeed does resolve that testcase, too. That wasn't obvious
>> to me. So, similar to
>> 'libgomp/testsuite/libgomp.oacc-c-c++-common/pr95270-{1.2}.c', please
>> include my new OpenACC/Fortran testcase (if that makes sense to you), and
>> reference PR95270 in the commit log.
>
> My new OpenACC/Fortran testcase got again broken ('libgomp: pointer
> target not mapped for attach') by Tobias' commit
> 102502e32ea4e8a75d6b252ba319d09d735d9aa7 "[OpenMP, Fortran] Add
> structure/derived-type element mapping",
> <http://mid.mail-archive.com/c5b43e02-d1d5-e7cf-c11c-6daf1e8f33c5@codesourcery.com>.
>
> Similar ('libgomp: attempt to attach null pointer') for your new
> 'libgomp.oacc-fortran/attach-descriptor-1.f90'.
>
> (Whether or not 'attach'ing 'NULL' should actually be allowed, is a
> separate topic for discussion.)
>
> So this patch here will (obviously) need to be adapted to what Tobias
> changed.
I see what you pushed in commit 39dda0020801045d9a604575b2a2593c05310015
"openacc: Fix standalone attach for Fortran assumed-shape array pointers"
indeed has become much smaller/simpler. :-)
(But, (parts of?) Tobias' commit mentioned above (plus commit
524862db444b6544c6dc87c5f06f351100ecf50d "Fix goacc/finalize-1.f tree
dump-scanning for -m32", if applicable) will then also need to be
backported to releases/gcc-10 branch (once un-frozen).)
> (Plus my more general questions quoted above and below.)
>>> OK?
>>
>> Basically yes (for master and releases/gcc-10 branches), but please
>> consider the following:
>>
>>> --- a/gcc/fortran/trans-openmp.c
>>> +++ b/gcc/fortran/trans-openmp.c
>>> @@ -2573,8 +2573,44 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
>>> }
>>> }
>>> if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (decl))
>>> - && n->u.map_op != OMP_MAP_ATTACH
>>> - && n->u.map_op != OMP_MAP_DETACH)
>>> + && (n->u.map_op == OMP_MAP_ATTACH
>>> + || n->u.map_op == OMP_MAP_DETACH))
>>> + {
>>> + tree type = TREE_TYPE (decl);
>>> + tree data = gfc_conv_descriptor_data_get (decl);
>>> + if (present)
>>> + data = gfc_build_cond_assign_expr (block, present,
>>> + data,
>>> + null_pointer_node);
>>> + tree ptr
>>> + = fold_convert (build_pointer_type (char_type_node),
>>> + data);
>>> + ptr = build_fold_indirect_ref (ptr);
>>> + /* Standalone attach clauses used with arrays with
>>> + descriptors must copy the descriptor to the target,
>>> + else they won't have anything to perform the
>>> + attachment onto (see OpenACC 2.6, "2.6.3. Data
>>> + Structures with Pointers"). */
>>> + OMP_CLAUSE_DECL (node) = ptr;
>>> + node2 = build_omp_clause (input_location, OMP_CLAUSE_MAP);
>>> + OMP_CLAUSE_SET_MAP_KIND (node2, GOMP_MAP_TO_PSET);
>>> + OMP_CLAUSE_DECL (node2) = decl;
>>> + OMP_CLAUSE_SIZE (node2) = TYPE_SIZE_UNIT (type);
>>> + node3 = build_omp_clause (input_location, OMP_CLAUSE_MAP);
>>> + if (n->u.map_op == OMP_MAP_ATTACH)
>>> + {
>>> + OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_ATTACH);
>>> + n->u.map_op = OMP_MAP_ALLOC;
>>> + }
>>> + else /* OMP_MAP_DETACH. */
>>> + {
>>> + OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_DETACH);
>>> + n->u.map_op = OMP_MAP_RELEASE;
>>> + }
>>> + OMP_CLAUSE_DECL (node3) = data;
>>> + OMP_CLAUSE_SIZE (node3) = size_int (0);
>>> + }
>>
>> So this ("case A") duplicates most of the code from...
>>
>>> + else if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (decl)))
>>> {
>>> [...]
>>
>> ... this existing case here ("case B"). It's not clear to me if these
>> two cases really still need to be handled separately, and a little bit
>> differently (regarding 'if (present)' handling, for example), or if they
>> could/should (?) be merged? Tobias, do you have an opinion?
(These have been merged.)
>> Do we have sufficient testsuite coverage? (For example,
>> 'attach'/'detach' with 'present == false', if that makes sense, or any
>> other thing that case A is doing differently from case B?)
(I'm not sure we're actually testing all relevant cases.)
>> Shouldn't
>> this get '-fdump-tree-original' and/or '-fdump-tree-gimple' testcases,
>> similar to 'gfortran.dg/goacc/finalize-1.f', so that we verify/document
>> what we generate here?
So I guess I had -- unconsciously? ;-) -- mentioned -fdump-tree-gimple'
and 'gfortran.dg/goacc/finalize-1.f' for a reason. That displays how the
'finalize' clause is implemented (see WIP patch attached,
'gfortran.dg/goacc/attach-descriptor.f90'), and...
>>> --- /dev/null
>>> +++ b/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90
>>> @@ -0,0 +1,51 @@
>>> +program att
>>
>> Please add 'dg-do run', and...
>>
>>> + use openacc
>>> + implicit none
>>> + type t
>>> + integer :: arr1(10)
>>> + integer, allocatable :: arr2(:)
>>> + end type t
>>> + integer :: i
>>> + type(t) :: myvar
>>> + integer, target :: tarr(10)
>>> + integer, pointer :: myptr(:)
>>> +
>>> + allocate(myvar%arr2(10))
>>> +
>>> + do i=1,10
>>> + myvar%arr1(i) = 0
>>> + myvar%arr2(i) = 0
>>> + tarr(i) = 0
>>> + end do
>>> +
>>> + call acc_copyin(myvar)
>>> + call acc_copyin(myvar%arr2)
>>> + call acc_copyin(tarr)
>>> +
>>> + myptr => tarr
>>> +
>>> + !$acc enter data attach(myvar%arr2, myptr)
>>> +
>>> + ! FIXME: This warning is emitted on the wrong line number.
>>> + ! { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } 36 }
>>
>> ... don't forget to adjust "36" here. ;-)
>>
>>> + !$acc serial present(myvar%arr2)
>>> + do i=1,10
>>> + myvar%arr1(i) = i
>>> + myvar%arr2(i) = i
>>> + end do
>>> + myptr(3) = 99
>>> + !$acc end serial
>>> +
>>> + !$acc exit data detach(myvar%arr2, myptr)
..., if we here 'detach' with 'finalize' added, that will turn into a
'delete' (instead of 'release') of 'myptr => tarr', and thus...
>>> +
>>> + call acc_copyout(myvar%arr2)
>>> + call acc_copyout(myvar)
>>> + call acc_copyout(tarr)
>>> +
>>> + do i=1,10
>>> + if (myvar%arr1(i) .ne. i) stop 1
>>> + if (myvar%arr2(i) .ne. i) stop 2
>>> + end do
>>> + if (tarr(3) .ne. 99) stop 3
..., here we won't see the updated 'tarr(3) == 99', and fail.
>>> +
>>> +end program att
Alternativly, we can show the problem with 'acc_is_present', as in my WIP
patch attached, 'libgomp.oacc-fortran/attach-descriptor-1__.f90'. (But
when experimenting with 'acc_is_present' and Fortran 'pointer's, beware
of PR96080 "OpenACC/Fortran runtime library routines vs. Fortran
'pointer'".)
What should happen in this case? Do we agree that 'exit data
detach(myptr)' should *never* unmap 'myptr => tarr', but really should
just unmap the 'myptr' array descriptor?
We can add special handling so that for standalone 'detach', a 'finalize'
doesn't turn 'release' into 'delete', but that doesn't feel like the
correct solution.
Also, we have a different -- bigger? -- problem: if we, for example,
'attach(myptr)' twice, that operation will include twice times
incrementing the reference count of 'myptr => tarr', and that'll then
conflict with a 'copyout(myptr)', as that one then sees unexpected
reference counts. That's a different variant of the "[OpenACC] Deep copy
attach/detach should not affect reference counts" problem?
Basically (see WIP patch attached,
'libgomp.oacc-fortran/attach-descriptor-1_.f90'):
call acc_copyin(tarr) ! 'rc(tarr) == 1'
myptr => tarr
!$acc enter data attach(myptr) ! 'rc(tarr) == 2'! (not intended by the user)
!$acc enter data attach(myptr) ! 'rc(tarr) == 3'! (not intended by the user)
[...]
call acc_copyout(tarr) ! won't copyout, because still 'rc(tarr) = 2'! (not intended by the user)
if (acc_is_present(tarr)) stop 12 ! fails
Ugh. :-( Or am I confused now?
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-WIP-Problems-with-openacc-Fix-standalone-attach-for-.patch --]
[-- Type: text/x-diff, Size: 7190 bytes --]
From 4fa4979da3de6d15d5a39b77fdeb6b5aadec0f10 Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <thomas@codesourcery.com>
Date: Fri, 17 Jul 2020 13:09:58 +0200
Subject: [PATCH] WIP Problems with "openacc: Fix standalone attach for Fortran
assumed-shape array pointers"
---
.../gfortran.dg/goacc/attach-descriptor.f90 | 11 ++++++++++-
.../libgomp.oacc-fortran/attach-descriptor-1.f90 | 3 +++
...descriptor-1.f90 => attach-descriptor-1_.f90} | 8 ++++++++
...escriptor-1.f90 => attach-descriptor-1__.f90} | 16 +++++++++++++++-
4 files changed, 36 insertions(+), 2 deletions(-)
copy libgomp/testsuite/libgomp.oacc-fortran/{attach-descriptor-1.f90 => attach-descriptor-1_.f90} (79%)
copy libgomp/testsuite/libgomp.oacc-fortran/{attach-descriptor-1.f90 => attach-descriptor-1__.f90} (62%)
diff --git a/gcc/testsuite/gfortran.dg/goacc/attach-descriptor.f90 b/gcc/testsuite/gfortran.dg/goacc/attach-descriptor.f90
index 9ca36f770c7..454ef9cccf3 100644
--- a/gcc/testsuite/gfortran.dg/goacc/attach-descriptor.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/attach-descriptor.f90
@@ -1,4 +1,4 @@
-! { dg-additional-options "-fdump-tree-original" }
+! { dg-additional-options "-fdump-tree-original -fdump-tree-gimple" }
program att
implicit none
@@ -12,7 +12,16 @@ program att
!$acc enter data attach(myvar%arr2, myptr)
! { dg-final { scan-tree-dump-times "(?n)#pragma acc enter data map\\(attach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(alloc:\\*\\(c_char \\*\\) myptr\\.data \\\[len: \[^\\\]\]+\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(attach:\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) myptr\\.data \\\[bias: 0\\\]\\);$" 1 "original" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(attach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(alloc:MEM\\\[\\(c_char \\*\\)\[^\\\]\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(attach:myptr\\.data \\\[bias: 0\\\]\\)$" 1 "gimple" } }
!$acc exit data detach(myvar%arr2, myptr)
! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(detach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(release:\\*\\(c_char \\*\\) myptr\\.data \\\[len: \[^\\\]\]+\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(detach:\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) myptr\\.data \\\[bias: 0\\\]\\);$" 1 "original" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(detach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(release:MEM\\\[\\(c_char \\*\\)\[^\\\]\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(detach:myptr\\.data \\\[bias: 0\\\]\\)$" 1 "gimple" } }
+
+ ! Test valid usage and processing of the finalize clause.
+ !$acc exit data detach(myvar%arr2, myptr) finalize
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(detach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(release:\\*\\(c_char \\*\\) myptr\\.data \\\[len: \[^\\\]\]+\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(detach:\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) myptr\\.data \\\[bias: 0\\\]\\) finalize;$" 1 "original" } }
+ !TODO See how in the 'gimple' dump, 'detach' is turned into 'force_detach', and 'release' into 'delete' -- but is the latter actually correct? (See 'libgomp.oacc-fortran/attach-descriptor-1__.f90'.)
+! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(force_detach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(delete:MEM\\\[\\(c_char \\*\\)\[^\\\]\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(force_detach:myptr\\.data \\\[bias: 0\\\]\\) finalize$" 1 "gimple" } }
+
end program att
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90
index 5d79cbc14fc..99c1e787de6 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90
@@ -41,8 +41,11 @@ program att
!$acc exit data detach(myvar%arr2, myptr)
call acc_copyout(myvar%arr2)
+ if (acc_is_present(myvar%arr2)) stop 10
call acc_copyout(myvar)
+ if (acc_is_present(myvar)) stop 11
call acc_copyout(tarr)
+ if (acc_is_present(tarr)) stop 12
do i=1,10
if (myvar%arr1(i) .ne. i) stop 1
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1_.f90
similarity index 79%
copy from libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90
copy to libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1_.f90
index 5d79cbc14fc..2e2e1267660 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1_.f90
@@ -26,6 +26,7 @@ program att
myptr => tarr
+ !$acc enter data attach(myvar%arr2, myptr)
!$acc enter data attach(myvar%arr2, myptr)
! FIXME: This warning is emitted on the wrong line number.
@@ -39,10 +40,17 @@ program att
!$acc end serial
!$acc exit data detach(myvar%arr2, myptr)
+ !!$acc exit data detach(myvar%arr2, myptr)
call acc_copyout(myvar%arr2)
+ if (acc_is_present(myvar%arr2)) stop 10
call acc_copyout(myvar)
+ if (acc_is_present(myvar)) stop 11
call acc_copyout(tarr)
+ if (acc_is_present(tarr)) stop 12 ! fails
+
+ !TODO Have to stop, have copied out device pointers.
+ stop
do i=1,10
if (myvar%arr1(i) .ne. i) stop 1
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1__.f90
similarity index 62%
copy from libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90
copy to libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1__.f90
index 5d79cbc14fc..6786f32852b 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1__.f90
@@ -27,6 +27,7 @@ program att
myptr => tarr
!$acc enter data attach(myvar%arr2, myptr)
+ !!$acc enter data attach(myvar%arr2, myptr)
! FIXME: This warning is emitted on the wrong line number.
! { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } 38 }
@@ -38,11 +39,24 @@ program att
myptr(3) = 99
!$acc end serial
- !$acc exit data detach(myvar%arr2, myptr)
+ !$acc exit data detach(myvar%arr2, myptr) finalize
+
+ if (.not. acc_is_present(myvar%arr2)) stop 10
+ if (.not. acc_is_present(myvar)) stop 11
+ if (.not. acc_is_present(tarr)) stop 12 ! fails
call acc_copyout(myvar%arr2)
+ if (acc_is_present(myvar%arr2)) stop 20
+ if (.not. acc_is_present(myvar)) stop 21
+ if (.not. acc_is_present(tarr)) stop 22
call acc_copyout(myvar)
+ if (acc_is_present(myvar%arr2)) stop 30
+ if (acc_is_present(myvar)) stop 31
+ if (.not. acc_is_present(tarr)) stop 32
call acc_copyout(tarr)
+ if (acc_is_present(myvar%arr2)) stop 40
+ if (acc_is_present(myvar)) stop 41
+ if (acc_is_present(tarr)) stop 42
do i=1,10
if (myvar%arr1(i) .ne. i) stop 1
--
2.17.1
^ permalink raw reply [flat|nested] 26+ messages in thread
* Re: [PATCH 9/9] [OpenACC] Don't detach for no-op exit data with zero dynamic refcount
2020-06-16 22:39 ` [PATCH 9/9] [OpenACC] Don't detach for no-op exit data with zero dynamic refcount Julian Brown
@ 2020-07-24 14:18 ` Thomas Schwinge
2020-07-24 22:53 ` Julian Brown
0 siblings, 1 reply; 26+ messages in thread
From: Thomas Schwinge @ 2020-07-24 14:18 UTC (permalink / raw)
To: Julian Brown
Cc: Jakub Jelinek, Tobias Burnus, Catherine_Moore, fortran, gcc-patches
Hi Julian!
On 2020-06-16T15:39:45-0700, Julian Brown <julian@codesourcery.com> wrote:
> This patch fixes a set of XFAILs
The overall goal of couse is not to resolve XFAILs, but to implement the
expected behavior. :-)
> in some recently-added patches by
> skipping a detach operation on "no-op" exit data operations for blocks
> with zero dynamic refcount.
So this is one aspect of <https://gcc.gnu.org/PR95203> "OpenACC 2.6 deep
copy: attach/detach API routines: no-op behavior".
> This takes advantage of the ordering of
> detach clauses with respect to associated data-movement clauses: i.e.,
> they are grouped together adjacently.
I'm not convinced that it's sufficient to just special-case these cases.
Instead, per the OpenACC "Data Clause Actions" etc., shouldn't basically
all 'gomp_fatal's that we have on 'attach'/'detach' code paths turn into
no-ops ("no action is taken")?
And I'd then again like to bring forward my idea from another review:
| we may (rather easily?) add a flag variable (ICV;
| initialized from an environment variable) to guard this checking
| behavior?
Here: to *keep* the 'gomp_fatal's.
| I suppose we may now have a few libgomp testcases that
| actually do [check things via expected 'gomp_fatal's],
| which wouldn't work any longer [...].
Such testcases could then 'dg-set-target-env-var "GOMP_ATTACH_FATAL" "1"'
(better name is desirable), and have one variant with and one variant
without that enabled.
Before you start re-working the patch, let's please first get agreement
on what exactly we intend to achieve.
Grüße
Thomas
> libgomp/
> * oacc-mem.c (find_group_last): Handle detach operations.
> (goacc_exit_data_internal): Detect detachments that are part of copyout
> operations, and suppress them if dynamic refcount is zero.
> * testsuite/libgomp.oacc-fortran/mdc-refcount-1-1-1.f90: Remove XFAILs.
> * testsuite/libgomp.oacc-fortran/mdc-refcount-1-1-2.F90: Fix typo.
> * testsuite/libgomp.oacc-fortran/mdc-refcount-1-2-1.f90: Remove XFAILs.
> * 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-3-2.f90: Likewise.
> * testsuite/libgomp.oacc-fortran/mdc-refcount-1-4-1.f90: Likewise.
> * testsuite/libgomp.oacc-fortran/mdc-refcount-1-4-2.f90: Likewise.
> ---
> libgomp/oacc-mem.c | 54 ++++++++++++++++---
> .../mdc-refcount-1-1-1.f90 | 6 +--
> .../mdc-refcount-1-1-2.F90 | 2 +-
> .../mdc-refcount-1-2-1.f90 | 6 +--
> .../mdc-refcount-1-2-2.f90 | 6 +--
> .../mdc-refcount-1-3-1.f90 | 6 +--
> .../mdc-refcount-1-3-2.f90 | 5 +-
> .../mdc-refcount-1-4-1.f90 | 6 +--
> .../mdc-refcount-1-4-2.f90 | 5 +-
> 9 files changed, 55 insertions(+), 41 deletions(-)
>
> diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c
> index 745cb132621..f852652c048 100644
> --- a/libgomp/oacc-mem.c
> +++ b/libgomp/oacc-mem.c
> @@ -987,7 +987,9 @@ find_group_last (int pos, size_t mapnum, size_t *sizes, unsigned short *kinds)
> {
> case GOMP_MAP_TO_PSET:
> if (pos + 1 < mapnum
> - && (kinds[pos + 1] & 0xff) == GOMP_MAP_ATTACH)
> + && ((kinds[pos + 1] & 0xff) == GOMP_MAP_ATTACH
> + || (kinds[pos + 1] & 0xff) == GOMP_MAP_DETACH
> + || (kinds[pos + 1] & 0xff) == GOMP_MAP_FORCE_DETACH))
> return pos + 1;
>
> while (pos + 1 < mapnum
> @@ -1010,6 +1012,8 @@ find_group_last (int pos, size_t mapnum, size_t *sizes, unsigned short *kinds)
> break;
>
> case GOMP_MAP_ATTACH:
> + case GOMP_MAP_DETACH:
> + case GOMP_MAP_FORCE_DETACH:
> return pos;
>
> default:
> @@ -1025,7 +1029,9 @@ find_group_last (int pos, size_t mapnum, size_t *sizes, unsigned short *kinds)
> /* We can have a single GOMP_MAP_ATTACH mapping after a to/from
> mapping. */
> if (pos + 1 < mapnum
> - && (kinds[pos + 1] & 0xff) == GOMP_MAP_ATTACH)
> + && ((kinds[pos + 1] & 0xff) == GOMP_MAP_ATTACH
> + || (kinds[pos + 1] & 0xff) == GOMP_MAP_DETACH
> + || (kinds[pos + 1] & 0xff) == GOMP_MAP_FORCE_DETACH))
> return pos + 1;
>
> /* We can have zero or more GOMP_MAP_POINTER mappings after a to/from
> @@ -1168,15 +1174,43 @@ goacc_exit_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum,
> {
> gomp_mutex_lock (&acc_dev->lock);
>
> - /* Handle "detach" before copyback/deletion of mapped data. */
> - for (size_t i = 0; i < mapnum; ++i)
> + /* Handle "detach" before copyback/deletion of mapped data. If this isn't a
> + standalone "detach" clause, take care to skip the "detach" operation if
> + the dynamic refcount of the data to be detached is zero. */
> + for (size_t grp = 0; grp < mapnum; grp++)
> {
> - unsigned char kind = kinds[i] & 0xff;
> + size_t i = grp, group_last = find_group_last (grp, mapnum, sizes, kinds);
> + unsigned char kind = kinds[grp] & 0xff;
> bool finalize = false;
> +
> switch (kind)
> {
> + case GOMP_MAP_TO_PSET:
> + case GOMP_MAP_TOFROM:
> + case GOMP_MAP_FROM:
> + case GOMP_MAP_FORCE_FROM:
> + case GOMP_MAP_RELEASE:
> + case GOMP_MAP_DELETE:
> + {
> + if (i + 1 >= mapnum)
> + break;
> + kind = kinds[i + 1] & 0xff;
> + if (kind != GOMP_MAP_FORCE_DETACH && kind != GOMP_MAP_DETACH)
> + break;
> + splay_tree_key n = lookup_host (acc_dev, hostaddrs[i], sizes[i]);
> + if (n == NULL)
> + {
> + gomp_mutex_unlock (&acc_dev->lock);
> + gomp_fatal ("target data not mapped for detach operation");
> + }
> + i++;
> + if (n->dynamic_refcount == 0)
> + break;
> + }
> + /* Fallthrough. */
> +
> case GOMP_MAP_FORCE_DETACH:
> - finalize = true;
> + finalize = (kind == GOMP_MAP_FORCE_DETACH);
> /* Fallthrough. */
>
> case GOMP_MAP_DETACH:
> @@ -1197,9 +1231,15 @@ goacc_exit_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum,
> gomp_detach_pointer (acc_dev, aq, n, hostaddr, finalize, NULL);
> }
> break;
> + case GOMP_MAP_STRUCT:
> + case GOMP_MAP_POINTER:
> + /* Ignore. */
> + break;
> default:
> - ;
> + gomp_fatal (">>>> goacc_exit_data_internal UNHANDLED kind 0x%.2x",
> + kind);
> }
> + grp = group_last;
> }
>
> for (size_t i = 0; i < mapnum; ++i)
> 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..7171affb9f0 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
> @@ -24,12 +24,8 @@ program main
> 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 } } }
> + ! { dg-output ".*CheCKpOInT2(\n|\r\n|\r)" }
> 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..2aa46189e9a 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
> @@ -6,4 +6,4 @@
> #include "mdc-refcount-1-1-1.f90"
>
> ! { dg-output ".*CheCKpOInT1(\n|\r\n|\r)" }
> -! { dg-output ".CheCKpOInT2(\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..9a10aa5a781 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
> @@ -26,12 +26,8 @@ program main
> 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 } } }
> + ! { dg-output ".*CheCKpOInT2(\n|\r\n|\r)" }
> 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..f506adf8e91 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
> @@ -26,12 +26,8 @@ program main
> 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 } } }
> + ! { dg-output ".*CheCKpOInT2(\n|\r\n|\r)" }
> 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..450d95d3686 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
> @@ -27,12 +27,8 @@ program main
> 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 } } }
> + ! { dg-output ".*CheCKpOInT2(\n|\r\n|\r)" }
> 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-2.f90 b/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-3-2.f90
> index 3c4bbda7f66..35efad4138a 100644
> --- a/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-3-2.f90
> +++ b/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-3-2.f90
> @@ -27,11 +27,8 @@ program main
> print *, "CheCKpOInT1"
> ! { dg-output ".*CheCKpOInT1(\n|\r\n|\r)" }
> !$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 } } }
> + ! { dg-output ".*CheCKpOInT2(\n|\r\n|\r)" }
> 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..816562fc055 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
> @@ -26,12 +26,8 @@ program main
> 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).
> - !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 } } }
> + ! { dg-output ".*CheCKpOInT2(\n|\r\n|\r)" }
> !$acc exit data delete(var%a)
> 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-2.f90 b/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-4-2.f90
> index 476cd5c1bee..b98bfd74924 100644
> --- a/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-4-2.f90
> +++ b/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-4-2.f90
> @@ -27,11 +27,8 @@ program main
> print *, "CheCKpOInT1"
> ! { dg-output ".*CheCKpOInT1(\n|\r\n|\r)" }
> !$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 } } }
> + ! { dg-output ".*CheCKpOInT2(\n|\r\n|\r)" }
> if (acc_is_present(var%a)) stop 3
> if (.not. acc_is_present(var)) stop 4
>
> --
> 2.23.0
-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter
^ permalink raw reply [flat|nested] 26+ messages in thread
* Re: [PATCH 9/9] [OpenACC] Don't detach for no-op exit data with zero dynamic refcount
2020-07-24 14:18 ` Thomas Schwinge
@ 2020-07-24 22:53 ` Julian Brown
0 siblings, 0 replies; 26+ messages in thread
From: Julian Brown @ 2020-07-24 22:53 UTC (permalink / raw)
To: Thomas Schwinge
Cc: Jakub Jelinek, Tobias Burnus, Catherine_Moore, fortran, gcc-patches
On Fri, 24 Jul 2020 16:18:34 +0200
Thomas Schwinge <thomas@codesourcery.com> wrote:
> Hi Julian!
>
> On 2020-06-16T15:39:45-0700, Julian Brown <julian@codesourcery.com>
> wrote:
> > This patch fixes a set of XFAILs
>
> The overall goal of couse is not to resolve XFAILs, but to implement
> the expected behavior. :-)
Eh, details :-).
> > in some recently-added patches by
> > skipping a detach operation on "no-op" exit data operations for
> > blocks with zero dynamic refcount.
>
> So this is one aspect of <https://gcc.gnu.org/PR95203> "OpenACC 2.6
> deep copy: attach/detach API routines: no-op behavior".
Not quite, I don't think -- that's about pointers (or pointed-to blocks)
that are not mapped on the target. With this patch, we have a detach
operation with a block with a zero dynamic refcount, but a non-zero
structured refcount -- i.e. it's still mapped. So I think the patch is
necessary, but not sufficient for the other cases you mention.
> > This takes advantage of the ordering of
> > detach clauses with respect to associated data-movement clauses:
> > i.e., they are grouped together adjacently.
>
> I'm not convinced that it's sufficient to just special-case these
> cases. Instead, per the OpenACC "Data Clause Actions" etc., shouldn't
> basically all 'gomp_fatal's that we have on 'attach'/'detach' code
> paths turn into no-ops ("no action is taken")?
>
>
> And I'd then again like to bring forward my idea from another review:
>
> | we may (rather easily?) add a flag variable (ICV;
> | initialized from an environment variable) to guard this checking
> | behavior?
>
> Here: to *keep* the 'gomp_fatal's.
>
> | I suppose we may now have a few libgomp testcases that
> | actually do [check things via expected 'gomp_fatal's],
> | which wouldn't work any longer [...].
>
> Such testcases could then 'dg-set-target-env-var "GOMP_ATTACH_FATAL"
> "1"' (better name is desirable), and have one variant with and one
> variant without that enabled.
>
>
> Before you start re-working the patch, let's please first get
> agreement on what exactly we intend to achieve.
Hm, you are probably right about the no-op behaviour for attach
operations (but slightly ugh in terms of usability), but I don't think
that's really the problem the patch addresses.
As for the user-tweakable checking -- yeah maybe it could work, but I
don't think I'm going to have time to work on that at the moment. Sorry!
HTH,
Julian
^ permalink raw reply [flat|nested] 26+ messages in thread
* Re: [PATCH 8/9] [OpenACC] Fix standalone attach for Fortran assumed-shape array pointers
2020-07-17 11:16 ` Thomas Schwinge
@ 2020-07-27 14:33 ` Julian Brown
2020-07-30 9:53 ` Thomas Schwinge
0 siblings, 1 reply; 26+ messages in thread
From: Julian Brown @ 2020-07-27 14:33 UTC (permalink / raw)
To: Thomas Schwinge
Cc: Tobias Burnus, gcc-patches, Jakub Jelinek, Catherine_Moore, fortran
[-- Attachment #1: Type: text/plain, Size: 8632 bytes --]
On Fri, 17 Jul 2020 13:16:11 +0200
Thomas Schwinge <thomas@codesourcery.com> wrote:
> Hi Julian, Tobias!
>
> On 2020-07-15T12:28:42+0200, Thomas Schwinge
> <thomas@codesourcery.com> wrote:
> > On 2020-07-14T13:43:37+0200, I wrote:
> >> On 2020-06-16T15:39:44-0700, Julian Brown
> >> <julian@codesourcery.com> wrote:
> >>> As mentioned in the blurb for the previous patch, an "attach"
> >>> operation for a Fortran pointer with an array descriptor must
> >>> copy that array descriptor to the target.
> >>
> >> Heh, I see -- I don't think I had read the OpenACC standard in
> >> that way, but I think I agree your interpretation is fine.
> >>
> >> This does not create some sort of memory leak -- everything
> >> implicitly allocated there will eventually be deallocated again,
> >> right?
>
> Unanswered -- but I may now have found this problem, and also found
> "the reverse problem" ('finalize'); see below.
Sorry, I didn't answer this explicitly -- the idea was to pair alloc
(present) and release mappings for the pointed-to data. In that way,
the idea was for the release mapping to perform that deallocation. That
was partly done so that the existing handling in gfc_trans_omp_clauses
could be used for this case without too much disruption to the code --
but actually, after Tobias's reorganisation of that function, that's
not really so much of an issue any more.
You can still get a "leak" if you try to attach a synthesized/temporary
array descriptor that goes out of scope before the pointed-to data it
refers to does -- that's a problem I've mentioned earlier, and is
kind-of unavoidable unless we do some more sophisticated analysis to
diagnose it as user error.
> >>> This patch arranges for that to be so.
> >>
> >> In response to the new OpenACC/Fortran testcase that I'd submtited
> >> in
> >> <http://mid.mail-archive.com/87wo3co0tm.fsf@euler.schwinge.homeip.net>,
> >> you (Julian) correctly supposed in
> >> <http://mid.mail-archive.com/20200709223246.23a4d0e0@squid.athome>,
> >> that this patch indeed does resolve that testcase, too. That
> >> wasn't obvious to me. So, similar to
> >> 'libgomp/testsuite/libgomp.oacc-c-c++-common/pr95270-{1.2}.c',
> >> please include my new OpenACC/Fortran testcase (if that makes
> >> sense to you), and reference PR95270 in the commit log.
> >
> > My new OpenACC/Fortran testcase got again broken ('libgomp: pointer
> > target not mapped for attach') by Tobias' commit
> > 102502e32ea4e8a75d6b252ba319d09d735d9aa7 "[OpenMP, Fortran] Add
> > structure/derived-type element mapping",
> > <http://mid.mail-archive.com/c5b43e02-d1d5-e7cf-c11c-6daf1e8f33c5@codesourcery.com>.
> >
> > Similar ('libgomp: attempt to attach null pointer') for your new
> > 'libgomp.oacc-fortran/attach-descriptor-1.f90'.
> >
> > (Whether or not 'attach'ing 'NULL' should actually be allowed, is a
> > separate topic for discussion.)
> >
> > So this patch here will (obviously) need to be adapted to what
> > Tobias changed.
>
> I see what you pushed in commit
> 39dda0020801045d9a604575b2a2593c05310015 "openacc: Fix standalone
> attach for Fortran assumed-shape array pointers" indeed has become
> much smaller/simpler. :-)
Yes, thank you.
> (But, (parts of?) Tobias' commit mentioned above (plus commit
> 524862db444b6544c6dc87c5f06f351100ecf50d "Fix goacc/finalize-1.f tree
> dump-scanning for -m32", if applicable) will then also need to be
> backported to releases/gcc-10 branch (once un-frozen).)
>
> > (Plus my more general questions quoted above and below.)
>
> >>> OK?
> >>
> >> Basically yes (for master and releases/gcc-10 branches), but please
> >> consider the following:
> >>
> >>> --- a/gcc/fortran/trans-openmp.c
> >>> +++ b/gcc/fortran/trans-openmp.c
> >>> @@ -2573,8 +2573,44 @@ gfc_trans_omp_clauses (stmtblock_t *block,
> >>> gfc_omp_clauses *clauses, }
> >>> }
> >>> if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (decl))
> >>> - && n->u.map_op != OMP_MAP_ATTACH
> >>> - && n->u.map_op != OMP_MAP_DETACH)
> >>> + && (n->u.map_op == OMP_MAP_ATTACH
> >>> + || n->u.map_op == OMP_MAP_DETACH))
> >>> + {
> >>> + tree type = TREE_TYPE (decl);
> >>> + tree data = gfc_conv_descriptor_data_get
> >>> (decl);
> >>> + if (present)
> >>> + data = gfc_build_cond_assign_expr
> >>> (block, present,
> >>> + data,
> >>> +
> >>> null_pointer_node);
> >>> + tree ptr
> >>> + = fold_convert (build_pointer_type
> >>> (char_type_node),
> >>> + data);
> >>> + ptr = build_fold_indirect_ref (ptr);
> >>> + /* Standalone attach clauses used with
> >>> arrays with
> >>> + descriptors must copy the descriptor to
> >>> the target,
> >>> + else they won't have anything to
> >>> perform the
> >>> + attachment onto (see OpenACC 2.6,
> >>> "2.6.3. Data
> >>> + Structures with Pointers"). */
> >>> + OMP_CLAUSE_DECL (node) = ptr;
> >>> + node2 = build_omp_clause (input_location,
> >>> OMP_CLAUSE_MAP);
> >>> + OMP_CLAUSE_SET_MAP_KIND (node2,
> >>> GOMP_MAP_TO_PSET);
> >>> + OMP_CLAUSE_DECL (node2) = decl;
> >>> + OMP_CLAUSE_SIZE (node2) = TYPE_SIZE_UNIT
> >>> (type);
> >>> + node3 = build_omp_clause (input_location,
> >>> OMP_CLAUSE_MAP);
> >>> + if (n->u.map_op == OMP_MAP_ATTACH)
> >>> + {
> >>> + OMP_CLAUSE_SET_MAP_KIND (node3,
> >>> GOMP_MAP_ATTACH);
> >>> + n->u.map_op = OMP_MAP_ALLOC;
> >>> + }
> >>> + else /* OMP_MAP_DETACH. */
> >>> + {
> >>> + OMP_CLAUSE_SET_MAP_KIND (node3,
> >>> GOMP_MAP_DETACH);
> >>> + n->u.map_op = OMP_MAP_RELEASE;
> >>> + }
> >>> + OMP_CLAUSE_DECL (node3) = data;
> >>> + OMP_CLAUSE_SIZE (node3) = size_int (0);
> >>> + }
> >>
> >> So this ("case A") duplicates most of the code from...
> >>
> >>> + else if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE
> >>> (decl))) {
> >>> [...]
> >>
> >> ... this existing case here ("case B"). It's not clear to me if
> >> these two cases really still need to be handled separately, and a
> >> little bit differently (regarding 'if (present)' handling, for
> >> example), or if they could/should (?) be merged? Tobias, do you
> >> have an opinion?
>
> (These have been merged.)
>
> >> Do we have sufficient testsuite coverage? (For example,
> >> 'attach'/'detach' with 'present == false', if that makes sense, or
> >> any other thing that case A is doing differently from case B?)
>
> (I'm not sure we're actually testing all relevant cases.)
...probably still not, sorry... more tests can be added later though of
course.
> >> Shouldn't
> >> this get '-fdump-tree-original' and/or '-fdump-tree-gimple'
> >> testcases, similar to 'gfortran.dg/goacc/finalize-1.f', so that we
> >> verify/document what we generate here?
>
> So I guess I had -- unconsciously? ;-) -- mentioned
> -fdump-tree-gimple' and 'gfortran.dg/goacc/finalize-1.f' for a
> reason. That displays how the 'finalize' clause is implemented (see
> WIP patch attached, 'gfortran.dg/goacc/attach-descriptor.f90'), and...
[snip]
> What should happen in this case? Do we agree that 'exit data
> detach(myptr)' should *never* unmap 'myptr => tarr', but really should
> just unmap the 'myptr' array descriptor?
>
> We can add special handling so that for standalone 'detach', a
> 'finalize' doesn't turn 'release' into 'delete', but that doesn't
> feel like the correct solution.
I don't think we actually need the alloc/release (with the latter turned
into "delete" for finalize) at all -- we just need to map the array
descriptor and perform the attach (or detach) as necessary. That's what
the attached patch does. Then, the pointed-to data's reference counts,
etc. will not be modified by attach/detach operations at all.
> Also, we have a different -- bigger? -- problem: if we, for example,
> 'attach(myptr)' twice, that operation will include twice times
> incrementing the reference count of 'myptr => tarr', and that'll then
> conflict with a 'copyout(myptr)', as that one then sees unexpected
> reference counts. That's a different variant of the "[OpenACC] Deep
> copy attach/detach should not affect reference counts" problem?
>
> Basically (see WIP patch attached,
> 'libgomp.oacc-fortran/attach-descriptor-1_.f90'):
Hmm, yes -- FWIW, this is caught by the "Refuse update/copyout for
blocks with attached pointers" patch. (In fact the attached patch
assumes that patch is already committed -- else the
attach-descriptor-4.f90 test should be XFAILed or omitted). So if we
want that one, this problem is sidestepped, I think.
Tested with offloading to NVPTX. OK?
Thanks,
Julian
[-- Attachment #2: 0001-openacc-No-attach-detach-present-release-mappings-fo.patch --]
[-- Type: text/x-patch, Size: 13487 bytes --]
From d53e4f1cd450062163e7e96a469c2f56cfac65ee Mon Sep 17 00:00:00 2001
From: Julian Brown <julian@codesourcery.com>
Date: Mon, 27 Jul 2020 06:29:02 -0700
Subject: [PATCH] openacc: No attach/detach present/release mappings for array
descriptors
Standalone attach and detach clauses should not create present/release
mappings for Fortran array descriptors (e.g. used when we have a pointer
to an array), both because it is unnecessary and because those mappings
will be incorrectly subject to reference counting. Simply omitting the
mappings means we just use GOMP_MAP_TO_PSET and GOMP_MAP_{ATTACH,DETACH}
mappings for array descriptors.
That requires a tweak in gimplify.c, since we may now see GOMP_MAP_TO_PSET
without a preceding data-movement mapping.
The new attach-descriptor-4.f90 test relies on the checking performed
by the patch "Refuse update/copyout for blocks with attached pointers".
2020-07-27 Julian Brown <julian@codesourcery.com>
Thomas Schwinge <thomas@codesourcery.com>
gcc/fortran/
* trans-openmp.c (gfc_trans_omp_clauses): Don't create present/release
mappings for array descriptors.
gcc/
* gimplify.c (gimplify_omp_target_update): Allow GOMP_MAP_TO_PSET
without a preceding data-movement mapping.
gcc/testsuite/
* gfortran.dg/goacc/attach-descriptor.f90: Update pattern output. Add
scanning of gimplify dump.
libgomp/
* testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90: Don't run for
shared-memory devices. Add more checking.
* testsuite/libgomp.oacc-fortran/attach-descriptor-3.f90: New test.
* testsuite/libgomp.oacc-fortran/attach-descriptor-4.f90: New test.
Co-Authored-By: Thomas Schwinge <thomas@codesourcery.com>
---
gcc/fortran/trans-openmp.c | 44 +++++++-----
gcc/gimplify.c | 3 +-
.../gfortran.dg/goacc/attach-descriptor.f90 | 17 ++++-
.../attach-descriptor-1.f90 | 6 +-
.../attach-descriptor-3.f90 | 68 +++++++++++++++++++
.../attach-descriptor-4.f90 | 61 +++++++++++++++++
6 files changed, 177 insertions(+), 22 deletions(-)
create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-3.f90
create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-4.f90
diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c
index d12d7fbddac..1a8f3277de3 100644
--- a/gcc/fortran/trans-openmp.c
+++ b/gcc/fortran/trans-openmp.c
@@ -2718,23 +2718,6 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
OMP_CLAUSE_SIZE (node2) = TYPE_SIZE_UNIT (type);
node3 = build_omp_clause (input_location,
OMP_CLAUSE_MAP);
- if (n->u.map_op == OMP_MAP_ATTACH)
- {
- /* Standalone attach clauses used with arrays with
- descriptors must copy the descriptor to the target,
- else they won't have anything to perform the
- attachment onto (see OpenACC 2.6, "2.6.3. Data
- Structures with Pointers"). */
- OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ALLOC);
- OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_ATTACH);
- }
- else if (n->u.map_op == OMP_MAP_DETACH)
- {
- OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_RELEASE);
- OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_DETACH);
- }
- else
- OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_POINTER);
if (present)
{
ptr = gfc_conv_descriptor_data_get (decl);
@@ -2748,6 +2731,33 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
OMP_CLAUSE_DECL (node3)
= gfc_conv_descriptor_data_get (decl);
OMP_CLAUSE_SIZE (node3) = size_int (0);
+ if (n->u.map_op == OMP_MAP_ATTACH)
+ {
+ /* Standalone attach clauses used with arrays with
+ descriptors must copy the descriptor to the target,
+ else they won't have anything to perform the
+ attachment onto (see OpenACC 2.6, "2.6.3. Data
+ Structures with Pointers"). */
+ OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_ATTACH);
+ /* We don't want to map PTR at all in this case, so
+ delete its node and shuffle the others down. */
+ node = node2;
+ node2 = node3;
+ node3 = NULL;
+ goto finalize_map_clause;
+ }
+ else if (n->u.map_op == OMP_MAP_DETACH)
+ {
+ OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_DETACH);
+ /* Similarly to above, we don't want to unmap PTR
+ here. */
+ node = node2;
+ node2 = node3;
+ node3 = NULL;
+ goto finalize_map_clause;
+ }
+ else
+ OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_POINTER);
/* We have to check for n->sym->attr.dimension because
of scalar coarrays. */
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 15dfee903ab..f4c31d2870d 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -13013,8 +13013,9 @@ gimplify_omp_target_update (tree *expr_p, gimple_seq *pre_p)
OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_DELETE);
have_clause = true;
break;
- case GOMP_MAP_POINTER:
case GOMP_MAP_TO_PSET:
+ break;
+ case GOMP_MAP_POINTER:
/* TODO PR92929: we may see these here, but they'll always follow
one of the clauses above, and will be handled by libgomp as
one group, so no handling required here. */
diff --git a/gcc/testsuite/gfortran.dg/goacc/attach-descriptor.f90 b/gcc/testsuite/gfortran.dg/goacc/attach-descriptor.f90
index 9ca36f770c7..373bdcb2114 100644
--- a/gcc/testsuite/gfortran.dg/goacc/attach-descriptor.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/attach-descriptor.f90
@@ -1,4 +1,4 @@
-! { dg-additional-options "-fdump-tree-original" }
+! { dg-additional-options "-fdump-tree-original -fdump-tree-gimple" }
program att
implicit none
@@ -11,8 +11,19 @@ program att
integer, pointer :: myptr(:)
!$acc enter data attach(myvar%arr2, myptr)
-! { dg-final { scan-tree-dump-times "(?n)#pragma acc enter data map\\(attach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(alloc:\\*\\(c_char \\*\\) myptr\\.data \\\[len: \[^\\\]\]+\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(attach:\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) myptr\\.data \\\[bias: 0\\\]\\);$" 1 "original" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc enter data map\\(attach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(attach:\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) myptr\\.data \\\[bias: 0\\\]\\);$" 1 "original" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(attach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(attach:myptr\\.data \\\[bias: 0\\\]\\)$" 1 "gimple" } }
!$acc exit data detach(myvar%arr2, myptr)
-! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(detach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(release:\\*\\(c_char \\*\\) myptr\\.data \\\[len: \[^\\\]\]+\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(detach:\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) myptr\\.data \\\[bias: 0\\\]\\);$" 1 "original" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(detach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(detach:\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) myptr\\.data \\\[bias: 0\\\]\\);$" 1 "original" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(detach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(detach:myptr\\.data \\\[bias: 0\\\]\\)$" 1 "gimple" } }
+
+ ! Test valid usage and processing of the finalize clause.
+ !$acc exit data detach(myvar%arr2, myptr) finalize
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(detach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(detach:\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) myptr\\.data \\\[bias: 0\\\]\\) finalize;$" 1 "original" } }
+ ! For array-descriptor detaches, we no longer generate a "release" mapping
+ ! for the pointed-to data for gimplify.c to turn into "delete". Make sure
+ ! the mapping still isn't there.
+! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(force_detach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(force_detach:myptr\\.data \\\[bias: 0\\\]\\) finalize$" 1 "gimple" } }
+
end program att
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90
index 5d79cbc14fc..9f159fa3b75 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90
@@ -1,4 +1,5 @@
! { dg-do run }
+! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } }
program att
use openacc
@@ -29,7 +30,7 @@ program att
!$acc enter data attach(myvar%arr2, myptr)
! FIXME: This warning is emitted on the wrong line number.
- ! { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } 38 }
+ ! { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } 39 }
!$acc serial present(myvar%arr2)
do i=1,10
myvar%arr1(i) = i
@@ -41,8 +42,11 @@ program att
!$acc exit data detach(myvar%arr2, myptr)
call acc_copyout(myvar%arr2)
+ if (acc_is_present(myvar%arr2)) stop 10
call acc_copyout(myvar)
+ if (acc_is_present(myvar)) stop 11
call acc_copyout(tarr)
+ if (acc_is_present(tarr)) stop 12
do i=1,10
if (myvar%arr1(i) .ne. i) stop 1
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-3.f90 b/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-3.f90
new file mode 100644
index 00000000000..f0e57b47453
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-3.f90
@@ -0,0 +1,68 @@
+! { dg-do run }
+! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } }
+
+program att
+ use openacc
+ implicit none
+ type t
+ integer :: arr1(10)
+ integer, allocatable :: arr2(:)
+ end type t
+ integer :: i
+ type(t) :: myvar
+ integer, target :: tarr(10)
+ integer, pointer :: myptr(:)
+
+ allocate(myvar%arr2(10))
+
+ do i=1,10
+ myvar%arr1(i) = 0
+ myvar%arr2(i) = 0
+ tarr(i) = 0
+ end do
+
+ call acc_copyin(myvar)
+ call acc_copyin(myvar%arr2)
+ call acc_copyin(tarr)
+
+ myptr => tarr
+
+ !$acc enter data attach(myvar%arr2, myptr)
+ !$acc enter data attach(myvar%arr2, myptr)
+
+ ! FIXME: This warning is emitted on the wrong line number.
+ ! { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } 40 }
+ !$acc serial present(myvar%arr2)
+ do i=1,10
+ myvar%arr1(i) = i
+ myvar%arr2(i) = i
+ end do
+ myptr(3) = 99
+ !$acc end serial
+
+ !$acc exit data detach(myvar%arr2, myptr) finalize
+
+ if (.not. acc_is_present(myvar%arr2)) stop 10
+ if (.not. acc_is_present(myvar)) stop 11
+ if (.not. acc_is_present(tarr)) stop 12
+
+ call acc_copyout(myvar%arr2)
+ if (acc_is_present(myvar%arr2)) stop 20
+ if (.not. acc_is_present(myvar)) stop 21
+ if (.not. acc_is_present(tarr)) stop 22
+ call acc_copyout(myvar)
+ if (acc_is_present(myvar%arr2)) stop 30
+ if (acc_is_present(myvar)) stop 31
+ if (.not. acc_is_present(tarr)) stop 32
+ call acc_copyout(tarr)
+ if (acc_is_present(myvar%arr2)) stop 40
+ if (acc_is_present(myvar)) stop 41
+ if (acc_is_present(tarr)) stop 42
+
+ do i=1,10
+ if (myvar%arr1(i) .ne. i) stop 1
+ if (myvar%arr2(i) .ne. i) stop 2
+ end do
+ if (tarr(3) .ne. 99) stop 3
+
+end program att
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-4.f90 b/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-4.f90
new file mode 100644
index 00000000000..9dbf53d0213
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-4.f90
@@ -0,0 +1,61 @@
+! { dg-do run }
+! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } }
+
+program att
+ use openacc
+ implicit none
+ type t
+ integer :: arr1(10)
+ integer, allocatable :: arr2(:)
+ end type t
+ integer :: i
+ type(t) :: myvar
+ integer, target :: tarr(10)
+ integer, pointer :: myptr(:)
+
+ allocate(myvar%arr2(10))
+
+ do i=1,10
+ myvar%arr1(i) = 0
+ myvar%arr2(i) = 0
+ tarr(i) = 0
+ end do
+
+ call acc_copyin(myvar)
+ call acc_copyin(myvar%arr2)
+ call acc_copyin(tarr)
+
+ myptr => tarr
+
+ !$acc enter data attach(myvar%arr2, myptr)
+ !$acc enter data attach(myvar%arr2, myptr)
+
+ ! FIXME: This warning is emitted on the wrong line number.
+ ! { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } 40 }
+ !$acc serial present(myvar%arr2)
+ do i=1,10
+ myvar%arr1(i) = i
+ myvar%arr2(i) = i
+ end do
+ myptr(3) = 99
+ !$acc end serial
+
+ !$acc exit data detach(myvar%arr2, myptr)
+
+ call acc_copyout(myvar%arr2)
+ ! { dg-output ".*copyout of block \\\[0x\[0-9a-f\]+,\\+\[0-9\]+\\\] with attached pointers(\n|\r\n|\r)+" }
+ if (acc_is_present(myvar%arr2)) stop 10
+ call acc_copyout(myvar)
+ if (acc_is_present(myvar)) stop 11
+ call acc_copyout(tarr)
+ if (acc_is_present(tarr)) stop 12
+
+ do i=1,10
+ if (myvar%arr1(i) .ne. i) stop 1
+ if (myvar%arr2(i) .ne. i) stop 2
+ end do
+ if (tarr(3) .ne. 99) stop 3
+
+end program att
+
+! { dg-shouldfail "" }
--
2.23.0
^ permalink raw reply [flat|nested] 26+ messages in thread
* Re: [PATCH 8/9] [OpenACC] Fix standalone attach for Fortran assumed-shape array pointers
2020-07-27 14:33 ` Julian Brown
@ 2020-07-30 9:53 ` Thomas Schwinge
2020-07-30 20:15 ` Julian Brown
0 siblings, 1 reply; 26+ messages in thread
From: Thomas Schwinge @ 2020-07-30 9:53 UTC (permalink / raw)
To: Julian Brown, Tobias Burnus
Cc: Jakub Jelinek, gcc-patches, fortran, Catherine_Moore
[-- Attachment #1: Type: text/plain, Size: 25339 bytes --]
Hi Julian, Tobias!
On 2020-07-27T15:33:41+0100, Julian Brown <julian@codesourcery.com> wrote:
> On Fri, 17 Jul 2020 13:16:11 +0200
> Thomas Schwinge <thomas@codesourcery.com> wrote:
>> On 2020-07-15T12:28:42+0200, Thomas Schwinge
>> <thomas@codesourcery.com> wrote:
>> > On 2020-07-14T13:43:37+0200, I wrote:
>> >> On 2020-06-16T15:39:44-0700, Julian Brown
>> >> <julian@codesourcery.com> wrote:
>> >>> As mentioned in the blurb for the previous patch, an "attach"
>> >>> operation for a Fortran pointer with an array descriptor must
>> >>> copy that array descriptor to the target.
>> >>
>> >> Heh, I see -- I don't think I had read the OpenACC standard in
>> >> that way, but I think I agree your interpretation is fine.
>> >>
>> >> This does not create some sort of memory leak -- everything
>> >> implicitly allocated there will eventually be deallocated again,
>> >> right?
>>
>> Unanswered -- but I may now have found this problem, and also found
>> "the reverse problem" ('finalize'); see below.
>
> Sorry, I didn't answer this explicitly -- the idea was to pair alloc
> (present) and release mappings for the pointed-to data. In that way,
> the idea was for the release mapping to perform that deallocation. That
> was partly done so that the existing handling in gfc_trans_omp_clauses
> could be used for this case without too much disruption to the code --
> but actually, after Tobias's reorganisation of that function, that's
> not really so much of an issue any more.
>
> You can still get a "leak" if you try to attach a synthesized/temporary
> array descriptor that goes out of scope before the pointed-to data it
> refers to does -- that's a problem I've mentioned earlier, and is
> kind-of unavoidable unless we do some more sophisticated analysis to
> diagnose it as user error.
ACK. Do you remember if you already had a testcase (conceptual, or
actual) to demonstrate that problem?
>> >>> This patch arranges for that to be so.
>> >>
>> >> In response to the new OpenACC/Fortran testcase that I'd submtited
>> >> in
>> >> <http://mid.mail-archive.com/87wo3co0tm.fsf@euler.schwinge.homeip.net>,
>> >> you (Julian) correctly supposed in
>> >> <http://mid.mail-archive.com/20200709223246.23a4d0e0@squid.athome>,
>> >> that this patch indeed does resolve that testcase, too. That
>> >> wasn't obvious to me. So, similar to
>> >> 'libgomp/testsuite/libgomp.oacc-c-c++-common/pr95270-{1.2}.c',
>> >> please include my new OpenACC/Fortran testcase (if that makes
>> >> sense to you), and reference PR95270 in the commit log.
>> >
>> > My new OpenACC/Fortran testcase got again broken ('libgomp: pointer
>> > target not mapped for attach') by Tobias' commit
>> > 102502e32ea4e8a75d6b252ba319d09d735d9aa7 "[OpenMP, Fortran] Add
>> > structure/derived-type element mapping",
>> > <http://mid.mail-archive.com/c5b43e02-d1d5-e7cf-c11c-6daf1e8f33c5@codesourcery.com>.
>> >
>> > Similar ('libgomp: attempt to attach null pointer') for your new
>> > 'libgomp.oacc-fortran/attach-descriptor-1.f90'.
>> >
>> > (Whether or not 'attach'ing 'NULL' should actually be allowed, is a
>> > separate topic for discussion.)
>> >
>> > So this patch here will (obviously) need to be adapted to what
>> > Tobias changed.
>>
>> I see what you pushed in commit
>> 39dda0020801045d9a604575b2a2593c05310015 "openacc: Fix standalone
>> attach for Fortran assumed-shape array pointers" indeed has become
>> much smaller/simpler. :-)
>
> Yes, thank you.
>
>> (But, (parts of?) Tobias' commit mentioned above (plus commit
>> 524862db444b6544c6dc87c5f06f351100ecf50d "Fix goacc/finalize-1.f tree
>> dump-scanning for -m32", if applicable) will then also need to be
>> backported to releases/gcc-10 branch (once un-frozen).)
>>
>> > (Plus my more general questions quoted above and below.)
>>
>> >>> OK?
>> >>
>> >> Basically yes (for master and releases/gcc-10 branches), but please
>> >> consider the following:
>> >>
>> >>> --- a/gcc/fortran/trans-openmp.c
>> >>> +++ b/gcc/fortran/trans-openmp.c
>> >>> @@ -2573,8 +2573,44 @@ gfc_trans_omp_clauses (stmtblock_t *block,
>> >>> gfc_omp_clauses *clauses, }
>> >>> }
>> >>> if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (decl))
>> >>> - && n->u.map_op != OMP_MAP_ATTACH
>> >>> - && n->u.map_op != OMP_MAP_DETACH)
>> >>> + && (n->u.map_op == OMP_MAP_ATTACH
>> >>> + || n->u.map_op == OMP_MAP_DETACH))
>> >>> + {
>> >>> + tree type = TREE_TYPE (decl);
>> >>> + tree data = gfc_conv_descriptor_data_get (decl);
>> >>> + if (present)
>> >>> + data = gfc_build_cond_assign_expr (block, present,
>> >>> + data,
>> >>> + null_pointer_node);
>> >>> + tree ptr
>> >>> + = fold_convert (build_pointer_type (char_type_node),
>> >>> + data);
>> >>> + ptr = build_fold_indirect_ref (ptr);
>> >>> + /* Standalone attach clauses used with arrays with
>> >>> + descriptors must copy the descriptor to the target,
>> >>> + else they won't have anything to perform the
>> >>> + attachment onto (see OpenACC 2.6, "2.6.3. Data
>> >>> + Structures with Pointers"). */
>> >>> + OMP_CLAUSE_DECL (node) = ptr;
>> >>> + node2 = build_omp_clause (input_location, OMP_CLAUSE_MAP);
>> >>> + OMP_CLAUSE_SET_MAP_KIND (node2, GOMP_MAP_TO_PSET);
>> >>> + OMP_CLAUSE_DECL (node2) = decl;
>> >>> + OMP_CLAUSE_SIZE (node2) = TYPE_SIZE_UNIT (type);
>> >>> + node3 = build_omp_clause (input_location, OMP_CLAUSE_MAP);
>> >>> + if (n->u.map_op == OMP_MAP_ATTACH)
>> >>> + {
>> >>> + OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_ATTACH);
>> >>> + n->u.map_op = OMP_MAP_ALLOC;
>> >>> + }
>> >>> + else /* OMP_MAP_DETACH. */
>> >>> + {
>> >>> + OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_DETACH);
>> >>> + n->u.map_op = OMP_MAP_RELEASE;
>> >>> + }
>> >>> + OMP_CLAUSE_DECL (node3) = data;
>> >>> + OMP_CLAUSE_SIZE (node3) = size_int (0);
>> >>> + }
>> >>
>> >> So this ("case A") duplicates most of the code from...
>> >>
>> >>> + else if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE
>> >>> (decl))) {
>> >>> [...]
>> >>
>> >> ... this existing case here ("case B"). It's not clear to me if
>> >> these two cases really still need to be handled separately, and a
>> >> little bit differently (regarding 'if (present)' handling, for
>> >> example), or if they could/should (?) be merged? Tobias, do you
>> >> have an opinion?
>>
>> (These have been merged.)
>>
>> >> Do we have sufficient testsuite coverage? (For example,
>> >> 'attach'/'detach' with 'present == false', if that makes sense, or
>> >> any other thing that case A is doing differently from case B?)
>>
>> (I'm not sure we're actually testing all relevant cases.)
>
> ...probably still not, sorry... more tests can be added later though of
> course.
(Just remains the question who's going to do that, "later"...) ;-\
>> >> Shouldn't
>> >> this get '-fdump-tree-original' and/or '-fdump-tree-gimple'
>> >> testcases, similar to 'gfortran.dg/goacc/finalize-1.f', so that we
>> >> verify/document what we generate here?
>>
>> So I guess I had -- unconsciously? ;-) -- mentioned
>> -fdump-tree-gimple' and 'gfortran.dg/goacc/finalize-1.f' for a
>> reason. That displays how the 'finalize' clause is implemented (see
>> WIP patch attached, 'gfortran.dg/goacc/attach-descriptor.f90'), and...
> [snip]
>> What should happen in this case? Do we agree that 'exit data
>> detach(myptr)' should *never* unmap 'myptr => tarr', but really should
>> just unmap the 'myptr' array descriptor?
>>
>> We can add special handling so that for standalone 'detach', a
>> 'finalize' doesn't turn 'release' into 'delete', but that doesn't
>> feel like the correct solution.
>
> I don't think we actually need the alloc/release (with the latter turned
> into "delete" for finalize) at all -- we just need to map the array
> descriptor and perform the attach (or detach) as necessary. That's what
> the attached patch does. Then, the pointed-to data's reference counts,
> etc. will not be modified by attach/detach operations at all.
ACK -- good to hear that this is the actual solution here.
>> Also, we have a different -- bigger? -- problem: if we, for example,
>> 'attach(myptr)' twice, that operation will include twice times
>> incrementing the reference count of 'myptr => tarr', and that'll then
>> conflict with a 'copyout(myptr)', as that one then sees unexpected
>> reference counts. That's a different variant of the "[OpenACC] Deep
>> copy attach/detach should not affect reference counts" problem?
>>
>> Basically (see WIP patch attached,
>> 'libgomp.oacc-fortran/attach-descriptor-1_.f90'):
>
> Hmm, yes -- FWIW, this is caught by the "Refuse update/copyout for
> blocks with attached pointers" patch. (In fact the attached patch
> assumes that patch is already committed -- else the
> attach-descriptor-4.f90 test should be XFAILed or omitted). So if we
> want that one, this problem is sidestepped, I think.
I'm attaching an incremental patch (I have tested that) to merge the
testcases into one file, and make it work on current master branch
without the pending "Refuse update/copyout for blocks with attached
pointers" changes. (We then later have to adjust the testcase here as
part of these changes.)
> Tested with offloading to NVPTX. OK?
Thanks. OK for master and releases/gcc-10 branches from my point of
view, but maybe Tobias can also have a look, please; two
comments/suggestions:
> From d53e4f1cd450062163e7e96a469c2f56cfac65ee Mon Sep 17 00:00:00 2001
> From: Julian Brown <julian@codesourcery.com>
> Date: Mon, 27 Jul 2020 06:29:02 -0700
> Subject: [PATCH] openacc: No attach/detach present/release mappings for array
> descriptors
>
> Standalone attach and detach clauses should not create present/release
> mappings for Fortran array descriptors (e.g. used when we have a pointer
> to an array), both because it is unnecessary and because those mappings
> will be incorrectly subject to reference counting. Simply omitting the
> mappings means we just use GOMP_MAP_TO_PSET and GOMP_MAP_{ATTACH,DETACH}
> mappings for array descriptors.
>
> That requires a tweak in gimplify.c, since we may now see GOMP_MAP_TO_PSET
> without a preceding data-movement mapping.
>
> The new attach-descriptor-4.f90 test relies on the checking performed
> by the patch "Refuse update/copyout for blocks with attached pointers".
(Need to remove that last sentence.)
> --- a/gcc/fortran/trans-openmp.c
> +++ b/gcc/fortran/trans-openmp.c
> @@ -2718,23 +2718,6 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
> OMP_CLAUSE_SIZE (node2) = TYPE_SIZE_UNIT (type);
> node3 = build_omp_clause (input_location,
> OMP_CLAUSE_MAP);
> - if (n->u.map_op == OMP_MAP_ATTACH)
> - {
> - /* Standalone attach clauses used with arrays with
> - descriptors must copy the descriptor to the target,
> - else they won't have anything to perform the
> - attachment onto (see OpenACC 2.6, "2.6.3. Data
> - Structures with Pointers"). */
> - OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ALLOC);
> - OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_ATTACH);
> - }
> - else if (n->u.map_op == OMP_MAP_DETACH)
> - {
> - OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_RELEASE);
> - OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_DETACH);
> - }
> - else
> - OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_POINTER);
> if (present)
> {
> ptr = gfc_conv_descriptor_data_get (decl);
> @@ -2748,6 +2731,33 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
> OMP_CLAUSE_DECL (node3)
> = gfc_conv_descriptor_data_get (decl);
> OMP_CLAUSE_SIZE (node3) = size_int (0);
> + if (n->u.map_op == OMP_MAP_ATTACH)
> + {
> + /* Standalone attach clauses used with arrays with
> + descriptors must copy the descriptor to the target,
> + else they won't have anything to perform the
> + attachment onto (see OpenACC 2.6, "2.6.3. Data
> + Structures with Pointers"). */
> + OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_ATTACH);
> + /* We don't want to map PTR at all in this case, so
> + delete its node and shuffle the others down. */
> + node = node2;
> + node2 = node3;
> + node3 = NULL;
> + goto finalize_map_clause;
> + }
> + else if (n->u.map_op == OMP_MAP_DETACH)
> + {
> + OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_DETACH);
> + /* Similarly to above, we don't want to unmap PTR
> + here. */
> + node = node2;
> + node2 = node3;
> + node3 = NULL;
> + goto finalize_map_clause;
> + }
> + else
> + OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_POINTER);
>
> /* We have to check for n->sym->attr.dimension because
> of scalar coarrays. */
I don't understand this code good enough to be sure that 'goto
finalize_map_clause' doesn't skip anything we may actually need -- for
the many "special" cases that Fortran has. Is it the case that it's the
correct thing to do, given that we're skipping 'node' completely.
I just had an idea how to make that clearer (maybe?) (untested, of
course): instead of the 'node', 'node2', 'node3' shuffling and 'goto
finalize_map_clause', don't do the shuffling and instead 'goto
finalize_map_clause_auxilliary' (better name maybe?):
finalize_map_clause:
omp_clauses = gfc_trans_add_clause (node, omp_clauses);
+finalize_map_clause_auxilliary:
if (node2)
omp_clauses = gfc_trans_add_clause (node2, omp_clauses);
if (node3)
omp_clauses = gfc_trans_add_clause (node3, omp_clauses);
if (node4)
omp_clauses = gfc_trans_add_clause (node4, omp_clauses);
(Just an idea; can also be done separately, later.)
> --- a/gcc/gimplify.c
> +++ b/gcc/gimplify.c
> @@ -13013,8 +13013,9 @@ gimplify_omp_target_update (tree *expr_p, gimple_seq *pre_p)
> OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_DELETE);
> have_clause = true;
> break;
> - case GOMP_MAP_POINTER:
> case GOMP_MAP_TO_PSET:
> + break;
> + case GOMP_MAP_POINTER:
> /* TODO PR92929: we may see these here, but they'll always follow
> one of the clauses above, and will be handled by libgomp as
> one group, so no handling required here. */
Maybe be good to add a comment why it's OK to do nothing for
'GOMP_MAP_TO_PSET'?
Grüße
Thomas
> --- a/gcc/testsuite/gfortran.dg/goacc/attach-descriptor.f90
> +++ b/gcc/testsuite/gfortran.dg/goacc/attach-descriptor.f90
> @@ -1,4 +1,4 @@
> -! { dg-additional-options "-fdump-tree-original" }
> +! { dg-additional-options "-fdump-tree-original -fdump-tree-gimple" }
>
> program att
> implicit none
> @@ -11,8 +11,19 @@ program att
> integer, pointer :: myptr(:)
>
> !$acc enter data attach(myvar%arr2, myptr)
> -! { dg-final { scan-tree-dump-times "(?n)#pragma acc enter data map\\(attach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(alloc:\\*\\(c_char \\*\\) myptr\\.data \\\[len: \[^\\\]\]+\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(attach:\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) myptr\\.data \\\[bias: 0\\\]\\);$" 1 "original" } }
> +! { dg-final { scan-tree-dump-times "(?n)#pragma acc enter data map\\(attach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(attach:\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) myptr\\.data \\\[bias: 0\\\]\\);$" 1 "original" } }
> +! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(attach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(attach:myptr\\.data \\\[bias: 0\\\]\\)$" 1 "gimple" } }
>
> !$acc exit data detach(myvar%arr2, myptr)
> -! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(detach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(release:\\*\\(c_char \\*\\) myptr\\.data \\\[len: \[^\\\]\]+\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(detach:\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) myptr\\.data \\\[bias: 0\\\]\\);$" 1 "original" } }
> +! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(detach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(detach:\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) myptr\\.data \\\[bias: 0\\\]\\);$" 1 "original" } }
> +! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(detach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(detach:myptr\\.data \\\[bias: 0\\\]\\)$" 1 "gimple" } }
> +
> + ! Test valid usage and processing of the finalize clause.
> + !$acc exit data detach(myvar%arr2, myptr) finalize
> +! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(detach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(detach:\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) myptr\\.data \\\[bias: 0\\\]\\) finalize;$" 1 "original" } }
> + ! For array-descriptor detaches, we no longer generate a "release" mapping
> + ! for the pointed-to data for gimplify.c to turn into "delete". Make sure
> + ! the mapping still isn't there.
> +! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(force_detach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(force_detach:myptr\\.data \\\[bias: 0\\\]\\) finalize$" 1 "gimple" } }
> +
> end program att
> diff --git a/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90
> index 5d79cbc14fc..9f159fa3b75 100644
> --- a/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90
> +++ b/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90
> @@ -1,4 +1,5 @@
> ! { dg-do run }
> +! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } }
>
> program att
> use openacc
> @@ -29,7 +30,7 @@ program att
> !$acc enter data attach(myvar%arr2, myptr)
>
> ! FIXME: This warning is emitted on the wrong line number.
> - ! { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } 38 }
> + ! { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } 39 }
> !$acc serial present(myvar%arr2)
> do i=1,10
> myvar%arr1(i) = i
> @@ -41,8 +42,11 @@ program att
> !$acc exit data detach(myvar%arr2, myptr)
>
> call acc_copyout(myvar%arr2)
> + if (acc_is_present(myvar%arr2)) stop 10
> call acc_copyout(myvar)
> + if (acc_is_present(myvar)) stop 11
> call acc_copyout(tarr)
> + if (acc_is_present(tarr)) stop 12
>
> do i=1,10
> if (myvar%arr1(i) .ne. i) stop 1
> diff --git a/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-3.f90 b/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-3.f90
> new file mode 100644
> index 00000000000..f0e57b47453
> --- /dev/null
> +++ b/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-3.f90
> @@ -0,0 +1,68 @@
> +! { dg-do run }
> +! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } }
> +
> +program att
> + use openacc
> + implicit none
> + type t
> + integer :: arr1(10)
> + integer, allocatable :: arr2(:)
> + end type t
> + integer :: i
> + type(t) :: myvar
> + integer, target :: tarr(10)
> + integer, pointer :: myptr(:)
> +
> + allocate(myvar%arr2(10))
> +
> + do i=1,10
> + myvar%arr1(i) = 0
> + myvar%arr2(i) = 0
> + tarr(i) = 0
> + end do
> +
> + call acc_copyin(myvar)
> + call acc_copyin(myvar%arr2)
> + call acc_copyin(tarr)
> +
> + myptr => tarr
> +
> + !$acc enter data attach(myvar%arr2, myptr)
> + !$acc enter data attach(myvar%arr2, myptr)
> +
> + ! FIXME: This warning is emitted on the wrong line number.
> + ! { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } 40 }
> + !$acc serial present(myvar%arr2)
> + do i=1,10
> + myvar%arr1(i) = i
> + myvar%arr2(i) = i
> + end do
> + myptr(3) = 99
> + !$acc end serial
> +
> + !$acc exit data detach(myvar%arr2, myptr) finalize
> +
> + if (.not. acc_is_present(myvar%arr2)) stop 10
> + if (.not. acc_is_present(myvar)) stop 11
> + if (.not. acc_is_present(tarr)) stop 12
> +
> + call acc_copyout(myvar%arr2)
> + if (acc_is_present(myvar%arr2)) stop 20
> + if (.not. acc_is_present(myvar)) stop 21
> + if (.not. acc_is_present(tarr)) stop 22
> + call acc_copyout(myvar)
> + if (acc_is_present(myvar%arr2)) stop 30
> + if (acc_is_present(myvar)) stop 31
> + if (.not. acc_is_present(tarr)) stop 32
> + call acc_copyout(tarr)
> + if (acc_is_present(myvar%arr2)) stop 40
> + if (acc_is_present(myvar)) stop 41
> + if (acc_is_present(tarr)) stop 42
> +
> + do i=1,10
> + if (myvar%arr1(i) .ne. i) stop 1
> + if (myvar%arr2(i) .ne. i) stop 2
> + end do
> + if (tarr(3) .ne. 99) stop 3
> +
> +end program att
> diff --git a/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-4.f90 b/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-4.f90
> new file mode 100644
> index 00000000000..9dbf53d0213
> --- /dev/null
> +++ b/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-4.f90
> @@ -0,0 +1,61 @@
> +! { dg-do run }
> +! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } }
> +
> +program att
> + use openacc
> + implicit none
> + type t
> + integer :: arr1(10)
> + integer, allocatable :: arr2(:)
> + end type t
> + integer :: i
> + type(t) :: myvar
> + integer, target :: tarr(10)
> + integer, pointer :: myptr(:)
> +
> + allocate(myvar%arr2(10))
> +
> + do i=1,10
> + myvar%arr1(i) = 0
> + myvar%arr2(i) = 0
> + tarr(i) = 0
> + end do
> +
> + call acc_copyin(myvar)
> + call acc_copyin(myvar%arr2)
> + call acc_copyin(tarr)
> +
> + myptr => tarr
> +
> + !$acc enter data attach(myvar%arr2, myptr)
> + !$acc enter data attach(myvar%arr2, myptr)
> +
> + ! FIXME: This warning is emitted on the wrong line number.
> + ! { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } 40 }
> + !$acc serial present(myvar%arr2)
> + do i=1,10
> + myvar%arr1(i) = i
> + myvar%arr2(i) = i
> + end do
> + myptr(3) = 99
> + !$acc end serial
> +
> + !$acc exit data detach(myvar%arr2, myptr)
> +
> + call acc_copyout(myvar%arr2)
> + ! { dg-output ".*copyout of block \\\[0x\[0-9a-f\]+,\\+\[0-9\]+\\\] with attached pointers(\n|\r\n|\r)+" }
> + if (acc_is_present(myvar%arr2)) stop 10
> + call acc_copyout(myvar)
> + if (acc_is_present(myvar)) stop 11
> + call acc_copyout(tarr)
> + if (acc_is_present(tarr)) stop 12
> +
> + do i=1,10
> + if (myvar%arr1(i) .ne. i) stop 1
> + if (myvar%arr2(i) .ne. i) stop 2
> + end do
> + if (tarr(3) .ne. 99) stop 3
> +
> +end program att
> +
> +! { dg-shouldfail "" }
-----------------
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-into-openacc-No-attach-detach-present-release-mappin.patch --]
[-- Type: text/x-diff, Size: 8293 bytes --]
From e3241486f68c077006513ea41c59ba3fdaeca7f7 Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <thomas@codesourcery.com>
Date: Wed, 29 Jul 2020 15:57:17 +0200
Subject: [PATCH] into: openacc: No attach/detach present/release mappings for
array descriptors
---
.../attach-descriptor-1.f90 | 93 ++++++++++++++++---
.../attach-descriptor-3.f90 | 68 --------------
.../attach-descriptor-4.f90 | 61 ------------
3 files changed, 80 insertions(+), 142 deletions(-)
delete mode 100644 libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-3.f90
delete mode 100644 libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-4.f90
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90
index 9f159fa3b75..960b9f94507 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90
@@ -1,9 +1,10 @@
! { dg-do run }
! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } }
-program att
+subroutine test(variant)
use openacc
implicit none
+ integer :: variant
type t
integer :: arr1(10)
integer, allocatable :: arr2(:)
@@ -27,31 +28,97 @@ program att
myptr => tarr
- !$acc enter data attach(myvar%arr2, myptr)
+ if (variant == 0 &
+ .or. variant == 3 &
+ .or. variant == 5) then
+ !$acc enter data attach(myvar%arr2, myptr)
+ else if (variant == 1 &
+ .or. variant == 2 &
+ .or. variant == 4) then
+ !$acc enter data attach(myvar%arr2, myptr)
+ !$acc enter data attach(myvar%arr2, myptr)
+ else
+ ! Internal error.
+ stop 1
+ end if
! FIXME: This warning is emitted on the wrong line number.
- ! { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } 39 }
+ ! { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } 52 }
!$acc serial present(myvar%arr2)
do i=1,10
- myvar%arr1(i) = i
- myvar%arr2(i) = i
+ myvar%arr1(i) = i + variant
+ myvar%arr2(i) = i - variant
end do
- myptr(3) = 99
+ myptr(3) = 99 - variant
!$acc end serial
- !$acc exit data detach(myvar%arr2, myptr)
+ if (variant == 0) then
+ !$acc exit data detach(myvar%arr2, myptr)
+ else if (variant == 1) then
+ !$acc exit data detach(myvar%arr2, myptr)
+ !$acc exit data detach(myvar%arr2, myptr)
+ else if (variant == 2) then
+ !$acc exit data detach(myvar%arr2, myptr)
+ !$acc exit data detach(myvar%arr2, myptr) finalize
+ else if (variant == 3 &
+ .or. variant == 4) then
+ !$acc exit data detach(myvar%arr2, myptr) finalize
+ else if (variant == 5) then
+ ! Do not detach.
+ else
+ ! Internal error.
+ stop 2
+ end if
+
+ if (.not. acc_is_present(myvar%arr2)) stop 10
+ if (.not. acc_is_present(myvar)) stop 11
+ if (.not. acc_is_present(tarr)) stop 12
call acc_copyout(myvar%arr2)
- if (acc_is_present(myvar%arr2)) stop 10
+ if (acc_is_present(myvar%arr2)) stop 20
+ if (.not. acc_is_present(myvar)) stop 21
+ if (.not. acc_is_present(tarr)) stop 22
call acc_copyout(myvar)
- if (acc_is_present(myvar)) stop 11
+ if (acc_is_present(myvar%arr2)) stop 30
+ if (acc_is_present(myvar)) stop 31
+ if (.not. acc_is_present(tarr)) stop 32
call acc_copyout(tarr)
- if (acc_is_present(tarr)) stop 12
+ if (acc_is_present(myvar%arr2)) stop 40
+ if (acc_is_present(myvar)) stop 41
+ if (acc_is_present(tarr)) stop 42
do i=1,10
- if (myvar%arr1(i) .ne. i) stop 1
- if (myvar%arr2(i) .ne. i) stop 2
+ if (myvar%arr1(i) .ne. i + variant) stop 50
+ if (variant == 5) then
+ ! We have not detached, so have copyied out a device pointer, so cannot
+ ! access 'myvar%arr2' on the host.
+ else
+ if (myvar%arr2(i) .ne. i - variant) stop 51
+ end if
end do
- if (tarr(3) .ne. 99) stop 3
+ if (tarr(3) .ne. 99 - variant) stop 52
+
+ if (variant == 5) then
+ ! If not explicitly stopping here, we'd in the following try to deallocate
+ ! the device pointer on the host, SIGSEGV.
+ stop
+ end if
+end subroutine test
+
+program att
+ implicit none
+
+ call test(0)
+
+ call test(1)
+
+ call test(2)
+
+ call test(3)
+
+ call test(4)
+ call test(5)
+ ! Make sure that 'test(5)' has stopped the program.
+ stop 60
end program att
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-3.f90 b/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-3.f90
deleted file mode 100644
index f0e57b47453..00000000000
--- a/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-3.f90
+++ /dev/null
@@ -1,68 +0,0 @@
-! { dg-do run }
-! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } }
-
-program att
- use openacc
- implicit none
- type t
- integer :: arr1(10)
- integer, allocatable :: arr2(:)
- end type t
- integer :: i
- type(t) :: myvar
- integer, target :: tarr(10)
- integer, pointer :: myptr(:)
-
- allocate(myvar%arr2(10))
-
- do i=1,10
- myvar%arr1(i) = 0
- myvar%arr2(i) = 0
- tarr(i) = 0
- end do
-
- call acc_copyin(myvar)
- call acc_copyin(myvar%arr2)
- call acc_copyin(tarr)
-
- myptr => tarr
-
- !$acc enter data attach(myvar%arr2, myptr)
- !$acc enter data attach(myvar%arr2, myptr)
-
- ! FIXME: This warning is emitted on the wrong line number.
- ! { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } 40 }
- !$acc serial present(myvar%arr2)
- do i=1,10
- myvar%arr1(i) = i
- myvar%arr2(i) = i
- end do
- myptr(3) = 99
- !$acc end serial
-
- !$acc exit data detach(myvar%arr2, myptr) finalize
-
- if (.not. acc_is_present(myvar%arr2)) stop 10
- if (.not. acc_is_present(myvar)) stop 11
- if (.not. acc_is_present(tarr)) stop 12
-
- call acc_copyout(myvar%arr2)
- if (acc_is_present(myvar%arr2)) stop 20
- if (.not. acc_is_present(myvar)) stop 21
- if (.not. acc_is_present(tarr)) stop 22
- call acc_copyout(myvar)
- if (acc_is_present(myvar%arr2)) stop 30
- if (acc_is_present(myvar)) stop 31
- if (.not. acc_is_present(tarr)) stop 32
- call acc_copyout(tarr)
- if (acc_is_present(myvar%arr2)) stop 40
- if (acc_is_present(myvar)) stop 41
- if (acc_is_present(tarr)) stop 42
-
- do i=1,10
- if (myvar%arr1(i) .ne. i) stop 1
- if (myvar%arr2(i) .ne. i) stop 2
- end do
- if (tarr(3) .ne. 99) stop 3
-
-end program att
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-4.f90 b/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-4.f90
deleted file mode 100644
index 9dbf53d0213..00000000000
--- a/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-4.f90
+++ /dev/null
@@ -1,61 +0,0 @@
-! { dg-do run }
-! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } }
-
-program att
- use openacc
- implicit none
- type t
- integer :: arr1(10)
- integer, allocatable :: arr2(:)
- end type t
- integer :: i
- type(t) :: myvar
- integer, target :: tarr(10)
- integer, pointer :: myptr(:)
-
- allocate(myvar%arr2(10))
-
- do i=1,10
- myvar%arr1(i) = 0
- myvar%arr2(i) = 0
- tarr(i) = 0
- end do
-
- call acc_copyin(myvar)
- call acc_copyin(myvar%arr2)
- call acc_copyin(tarr)
-
- myptr => tarr
-
- !$acc enter data attach(myvar%arr2, myptr)
- !$acc enter data attach(myvar%arr2, myptr)
-
- ! FIXME: This warning is emitted on the wrong line number.
- ! { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } 40 }
- !$acc serial present(myvar%arr2)
- do i=1,10
- myvar%arr1(i) = i
- myvar%arr2(i) = i
- end do
- myptr(3) = 99
- !$acc end serial
-
- !$acc exit data detach(myvar%arr2, myptr)
-
- call acc_copyout(myvar%arr2)
- ! { dg-output ".*copyout of block \\\[0x\[0-9a-f\]+,\\+\[0-9\]+\\\] with attached pointers(\n|\r\n|\r)+" }
- if (acc_is_present(myvar%arr2)) stop 10
- call acc_copyout(myvar)
- if (acc_is_present(myvar)) stop 11
- call acc_copyout(tarr)
- if (acc_is_present(tarr)) stop 12
-
- do i=1,10
- if (myvar%arr1(i) .ne. i) stop 1
- if (myvar%arr2(i) .ne. i) stop 2
- end do
- if (tarr(3) .ne. 99) stop 3
-
-end program att
-
-! { dg-shouldfail "" }
--
2.17.1
^ permalink raw reply [flat|nested] 26+ messages in thread
* Re: [PATCH 8/9] [OpenACC] Fix standalone attach for Fortran assumed-shape array pointers
2020-07-30 9:53 ` Thomas Schwinge
@ 2020-07-30 20:15 ` Julian Brown
0 siblings, 0 replies; 26+ messages in thread
From: Julian Brown @ 2020-07-30 20:15 UTC (permalink / raw)
To: Thomas Schwinge
Cc: Tobias Burnus, Jakub Jelinek, gcc-patches, fortran, Catherine_Moore
[-- Attachment #1: Type: text/plain, Size: 802 bytes --]
On Thu, 30 Jul 2020 11:53:10 +0200
Thomas Schwinge <thomas@codesourcery.com> wrote:
> Hi Julian, Tobias!
>
> On 2020-07-27T15:33:41+0100, Julian Brown <julian@codesourcery.com>
> wrote:
> > You can still get a "leak" if you try to attach a
> > synthesized/temporary array descriptor that goes out of scope
> > before the pointed-to data it refers to does -- that's a problem
> > I've mentioned earlier, and is kind-of unavoidable unless we do
> > some more sophisticated analysis to diagnose it as user error.
>
> ACK. Do you remember if you already had a testcase (conceptual, or
> actual) to demonstrate that problem?
I have the attached, but it's not "clean", i.e. not really
testsuite-ready -- the breakage demonstrated depends on the stack
layout, and it only "works" at -O0.
HTH,
Julian
[-- Attachment #2: enter-data-pset.f90 --]
[-- Type: text/x-fortran, Size: 1442 bytes --]
program myprog
implicit none
integer :: a(16)
integer :: i
call enterdata_wrapper(a, 16)
call exitdata_wrapper(a, 16)
contains
subroutine enterdata_wrapper(a, n)
implicit none
integer :: n
integer, target :: a(n)
integer :: aa(16)
integer :: bb(16)
integer, pointer :: ap(:)
integer :: cc(16)
integer :: dd(16)
! An array descriptor appears somewhere around here...
ap => a
!$acc enter data copyin(aa,bb,cc,dd)
call enterdata(ap)
!$acc exit data copyout(aa,bb,cc,dd)
! ...and goes out of scope.
end subroutine enterdata_wrapper
subroutine enterdata(a)
implicit none
integer, pointer :: a(:)
! Map "to(a.data) to_pset(a) pointer(a.data)"
!$acc enter data copyin(a)
end subroutine enterdata
subroutine exitdata_wrapper(a, n)
implicit none
integer :: n
integer, target :: a(n)
integer :: aa(32)
integer :: bb(32)
integer, pointer :: ap(:)
integer :: cc(32)
integer :: dd(32)
! A different array descriptor appears...
ap => a
!$acc enter data copyin(aa,bb,cc,dd)
call exitdata(ap)
!$acc exit data copyout(aa,bb,cc,dd)
! ...and goes out of scope.
end subroutine exitdata_wrapper
subroutine exitdata(a)
implicit none
integer, pointer :: a(:)
! Try to unmap the fresh array descriptor: FAILS.
!$acc exit data copyout(a)
end subroutine exitdata
end program myprog
^ permalink raw reply [flat|nested] 26+ messages in thread
end of thread, other threads:[~2020-07-30 20:15 UTC | newest]
Thread overview: 26+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2020-06-16 22:38 [PATCH 0/9] [OpenACC] Refcounting and manual deep copy improvements Julian Brown
2020-06-16 22:38 ` [PATCH 1/9] [OpenACC] Fortran derived-type mapping fix Julian Brown
2020-06-16 22:38 ` [PATCH 2/9] [OpenACC] GOMP_MAP_ATTACH handling in find_group_last Julian Brown
2020-06-30 12:42 ` Thomas Schwinge
2020-06-16 22:38 ` [PATCH 3/9] [OpenACC] Adjust dynamic reference count semantics Julian Brown
2020-06-30 13:51 ` Thomas Schwinge
2020-07-03 15:41 ` Thomas Schwinge
2020-07-10 12:08 ` Julian Brown
2020-06-16 22:38 ` [PATCH 4/9] [OpenACC] Don't pass kind array via pointer to goacc_enter_datum Julian Brown
2020-06-16 22:39 ` [PATCH 5/9] [OpenACC] Fix incompatible copyout for acc_map_data (PR92843) Julian Brown
2020-06-16 22:39 ` [PATCH 6/9] [OpenACC] Set bias to zero for explicit attach/detach clauses in C and C++ Julian Brown
2020-06-25 11:36 ` Thomas Schwinge
2020-07-09 21:06 ` Thomas Schwinge
2020-07-09 21:32 ` Julian Brown
2020-06-16 22:39 ` [PATCH 7/9] [OpenACC] Do not strip GOMP_MAP_TO_PSET/GOMP_MAP_POINTER for enter/exit data directives Julian Brown
2020-07-06 16:19 ` Thomas Schwinge
2020-06-16 22:39 ` [PATCH 8/9] [OpenACC] Fix standalone attach for Fortran assumed-shape array pointers Julian Brown
2020-07-14 11:43 ` Thomas Schwinge
2020-07-15 10:28 ` Thomas Schwinge
2020-07-17 11:16 ` Thomas Schwinge
2020-07-27 14:33 ` Julian Brown
2020-07-30 9:53 ` Thomas Schwinge
2020-07-30 20:15 ` Julian Brown
2020-06-16 22:39 ` [PATCH 9/9] [OpenACC] Don't detach for no-op exit data with zero dynamic refcount Julian Brown
2020-07-24 14:18 ` Thomas Schwinge
2020-07-24 22:53 ` Julian Brown
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).