Hi!
On Fri, 25 May 2018 13:01:58 -0700, Cesar Philippidis <ce...@codesourcery.com>
wrote:
> This patch updates GCC's to support OpenACC 2.5's data clause semantics.
> In OpenACC 2.5, copy, copyin and copyout all behave like their
> present_or_* counterparts in OpenACC 2.0. The patch also adds support
> for the new finalize and if_present data clauses introduced in OpenACC
> 2.5. The finalize clause introduced some new reference counting behavior
> in the runtime; whereas 'acc exit data copyout' decrements the reference
> count of a variable, 'acc exit data finalize' actually removes it from
> the accelerator regardless of there are any pending references to it.
>
> Due to the size of this patch, I had to compress it.
Well, or just handled separately the pieces that already were nicely
separated, instead of combining them into one big patch?
> However, despite
> the size of the patch, which is mainly due to the augmented test cases,
> it fairly noninvasive. I was originally going to include support for
> declare allocate and deallocate, but those require more extensive
> modifications to the fortran FE.
(Why the idea of including even more independent changes?)
Anyway. Committed to trunk in r267153 the following to add in the
missing changes from Chung-Lin's "Adjust copy/copyin/copyout/create for
OpenACC 2.5" patch:
commit 8ccac5746ba73bde3a1db490359bbe2d5f62d4a0
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Fri Dec 14 20:43:12 2018 +0000
Missing changes from "Adjust copy/copyin/copyout/create for OpenACC 2.5"
Most of that patch's changes were already committed as part of r261813
"Update
OpenACC data clause semantics to the 2.5 behavior", but not all of them.
libgomp/
* oacc-mem.c (acc_present_or_create): Remove definition and change
to alias of acc_create.
(acc_present_or_copyin): Remove definition and change to alias of
acc_copyin.
* oacc-parallel.c (GOACC_enter_exit_data): Call acc_create instead
of acc_present_or_create.
* testsuite/libgomp.oacc-c-c++-common/data-already-1.c: Remove.
* testsuite/libgomp.oacc-c-c++-common/data-already-2.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/data-already-3.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/data-already-4.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/data-already-5.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/data-already-6.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/data-already-7.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/data-already-8.c: Likewise.
* testsuite/libgomp.oacc-fortran/data-already-1.f: Likewise.
* testsuite/libgomp.oacc-fortran/data-already-2.f: Likewise.
* testsuite/libgomp.oacc-fortran/data-already-3.f: Likewise.
* testsuite/libgomp.oacc-fortran/data-already-4.f: Likewise.
* testsuite/libgomp.oacc-fortran/data-already-5.f: Likewise.
* testsuite/libgomp.oacc-fortran/data-already-6.f: Likewise.
* testsuite/libgomp.oacc-fortran/data-already-7.f: Likewise.
* testsuite/libgomp.oacc-fortran/data-already-8.f: Likewise.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@267153
138bc75d-0d04-0410-961f-82ee72b054a4
---
libgomp/ChangeLog | 26 ++++++++++
libgomp/oacc-mem.c | 55 ++++++++++++----------
libgomp/oacc-parallel.c | 4 --
.../libgomp.oacc-c-c++-common/data-already-1.c | 20 --------
.../libgomp.oacc-c-c++-common/data-already-2.c | 20 --------
.../libgomp.oacc-c-c++-common/data-already-3.c | 20 --------
.../libgomp.oacc-c-c++-common/data-already-4.c | 18 -------
.../libgomp.oacc-c-c++-common/data-already-5.c | 18 -------
.../libgomp.oacc-c-c++-common/data-already-6.c | 18 -------
.../libgomp.oacc-c-c++-common/data-already-7.c | 18 -------
.../libgomp.oacc-c-c++-common/data-already-8.c | 20 --------
.../libgomp.oacc-fortran/data-already-1.f | 16 -------
.../libgomp.oacc-fortran/data-already-2.f | 16 -------
.../libgomp.oacc-fortran/data-already-3.f | 15 ------
.../libgomp.oacc-fortran/data-already-4.f | 14 ------
.../libgomp.oacc-fortran/data-already-5.f | 14 ------
.../libgomp.oacc-fortran/data-already-6.f | 14 ------
.../libgomp.oacc-fortran/data-already-7.f | 14 ------
.../libgomp.oacc-fortran/data-already-8.f | 16 -------
19 files changed, 55 insertions(+), 301 deletions(-)
diff --git libgomp/ChangeLog libgomp/ChangeLog
index 349497d58ee6..084f174513b0 100644
--- libgomp/ChangeLog
+++ libgomp/ChangeLog
@@ -1,3 +1,29 @@
+2018-12-14 Thomas Schwinge <tho...@codesourcery.com>
+ Chung-Lin Tang <clt...@codesourcery.com>
+
+ * oacc-mem.c (acc_present_or_create): Remove definition and change
+ to alias of acc_create.
+ (acc_present_or_copyin): Remove definition and change to alias of
+ acc_copyin.
+ * oacc-parallel.c (GOACC_enter_exit_data): Call acc_create instead
+ of acc_present_or_create.
+ * testsuite/libgomp.oacc-c-c++-common/data-already-1.c: Remove.
+ * testsuite/libgomp.oacc-c-c++-common/data-already-2.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/data-already-3.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/data-already-4.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/data-already-5.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/data-already-6.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/data-already-7.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/data-already-8.c: Likewise.
+ * testsuite/libgomp.oacc-fortran/data-already-1.f: Likewise.
+ * testsuite/libgomp.oacc-fortran/data-already-2.f: Likewise.
+ * testsuite/libgomp.oacc-fortran/data-already-3.f: Likewise.
+ * testsuite/libgomp.oacc-fortran/data-already-4.f: Likewise.
+ * testsuite/libgomp.oacc-fortran/data-already-5.f: Likewise.
+ * testsuite/libgomp.oacc-fortran/data-already-6.f: Likewise.
+ * testsuite/libgomp.oacc-fortran/data-already-7.f: Likewise.
+ * testsuite/libgomp.oacc-fortran/data-already-8.f: Likewise.
+
2018-12-14 Thomas Schwinge <tho...@codesourcery.com>
PR libgomp/88495
diff --git libgomp/oacc-mem.c libgomp/oacc-mem.c
index 72414b748971..e352379bfeb3 100644
--- libgomp/oacc-mem.c
+++ libgomp/oacc-mem.c
@@ -544,6 +544,25 @@ acc_create_async (void *h, size_t s, int async)
present_create_copy (FLAG_PRESENT | FLAG_CREATE, h, s, async);
}
+/* acc_present_or_create used to be what acc_create is now. */
+/* acc_pcreate is acc_present_or_create by a different name. */
+#ifdef HAVE_ATTRIBUTE_ALIAS
+strong_alias (acc_create, acc_present_or_create)
+strong_alias (acc_create, acc_pcreate)
+#else
+void *
+acc_present_or_create (void *h, size_t s)
+{
+ return acc_create (h, s);
+}
+
+void *
+acc_pcreate (void *h, size_t s)
+{
+ return acc_create (h, s);
+}
+#endif
+
void *
acc_copyin (void *h, size_t s)
{
@@ -557,38 +576,22 @@ acc_copyin_async (void *h, size_t s, int async)
present_create_copy (FLAG_PRESENT | FLAG_CREATE | FLAG_COPY, h, s, async);
}
-void *
-acc_present_or_create (void *h, size_t s)
-{
- return present_create_copy (FLAG_PRESENT | FLAG_CREATE, h, s,
acc_async_sync);
-}
-
-/* acc_pcreate is acc_present_or_create by a different name. */
-#ifdef HAVE_ATTRIBUTE_ALIAS
-strong_alias (acc_present_or_create, acc_pcreate)
-#else
-void *
-acc_pcreate (void *h, size_t s)
-{
- return acc_present_or_create (h, s);
-}
-#endif
-
-void *
-acc_present_or_copyin (void *h, size_t s)
-{
- return present_create_copy (FLAG_PRESENT | FLAG_CREATE | FLAG_COPY, h, s,
- acc_async_sync);
-}
-
+/* acc_present_or_copyin used to be what acc_copyin is now. */
/* acc_pcopyin is acc_present_or_copyin by a different name. */
#ifdef HAVE_ATTRIBUTE_ALIAS
-strong_alias (acc_present_or_copyin, acc_pcopyin)
+strong_alias (acc_copyin, acc_present_or_copyin)
+strong_alias (acc_copyin, acc_pcopyin)
#else
+void *
+acc_present_or_copyin (void *h, size_t s)
+{
+ return acc_copyin (h, s);
+}
+
void *
acc_pcopyin (void *h, size_t s)
{
- return acc_present_or_copyin (h, s);
+ return acc_copyin (h, s);
}
#endif
diff --git libgomp/oacc-parallel.c libgomp/oacc-parallel.c
index 89b6b6f6fc2b..e2405d73d5a9 100644
--- libgomp/oacc-parallel.c
+++ libgomp/oacc-parallel.c
@@ -425,14 +425,10 @@ GOACC_enter_exit_data (int device, size_t mapnum,
switch (kind)
{
case GOMP_MAP_ALLOC:
- acc_present_or_create (hostaddrs[i], sizes[i]);
- break;
case GOMP_MAP_FORCE_ALLOC:
acc_create (hostaddrs[i], sizes[i]);
break;
case GOMP_MAP_TO:
- acc_present_or_copyin (hostaddrs[i], sizes[i]);
- break;
case GOMP_MAP_FORCE_TO:
acc_copyin (hostaddrs[i], sizes[i]);
break;
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-1.c
libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-1.c
deleted file mode 100644
index fd3b77dcff5a..000000000000
--- libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-1.c
+++ /dev/null
@@ -1,20 +0,0 @@
-/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
-
-#include <stdio.h>
-#include <openacc.h>
-
-int
-main (int argc, char *argv[])
-{
- int i;
-
- acc_copyin (&i, sizeof i);
-
- fprintf (stderr, "CheCKpOInT\n");
-#pragma acc data copy (i)
- ++i;
-
- return 0;
-}
-
-/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-2.c
libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-2.c
deleted file mode 100644
index 0118b2568e23..000000000000
--- libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-2.c
+++ /dev/null
@@ -1,20 +0,0 @@
-/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
-
-#include <stdio.h>
-
-int
-main (int argc, char *argv[])
-{
- int i;
-
-#pragma acc data present_or_copy (i)
- {
- fprintf (stderr, "CheCKpOInT\n");
-#pragma acc data copyout (i)
- ++i;
- }
-
- return 0;
-}
-
-/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-3.c
libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-3.c
deleted file mode 100644
index b346c69826fb..000000000000
--- libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-3.c
+++ /dev/null
@@ -1,20 +0,0 @@
-/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
-
-#include <stdio.h>
-#include <openacc.h>
-
-int
-main (int argc, char *argv[])
-{
- int i;
-
-#pragma acc data present_or_copy (i)
- {
- fprintf (stderr, "CheCKpOInT\n");
- acc_copyin (&i, sizeof i);
- }
-
- return 0;
-}
-
-/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-4.c
libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-4.c
deleted file mode 100644
index e99ad33d9bea..000000000000
--- libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-4.c
+++ /dev/null
@@ -1,18 +0,0 @@
-/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
-
-#include <stdio.h>
-#include <openacc.h>
-
-int
-main (int argc, char *argv[])
-{
- int i;
-
- acc_present_or_copyin (&i, sizeof i);
- fprintf (stderr, "CheCKpOInT\n");
- acc_copyin (&i, sizeof i);
-
- return 0;
-}
-
-/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-5.c
libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-5.c
deleted file mode 100644
index f8370c006df1..000000000000
--- libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-5.c
+++ /dev/null
@@ -1,18 +0,0 @@
-/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
-
-#include <stdio.h>
-#include <openacc.h>
-
-int
-main (int argc, char *argv[])
-{
- int i;
-
-#pragma acc enter data create (i)
- fprintf (stderr, "CheCKpOInT\n");
- acc_copyin (&i, sizeof i);
-
- return 0;
-}
-
-/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-6.c
libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-6.c
deleted file mode 100644
index d7f4deb18e48..000000000000
--- libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-6.c
+++ /dev/null
@@ -1,18 +0,0 @@
-/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
-
-#include <stdio.h>
-#include <openacc.h>
-
-int
-main (int argc, char *argv[])
-{
- int i;
-
- acc_present_or_copyin (&i, sizeof i);
- fprintf (stderr, "CheCKpOInT\n");
-#pragma acc enter data create (i)
-
- return 0;
-}
-
-/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-7.c
libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-7.c
deleted file mode 100644
index 54be59507ca0..000000000000
--- libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-7.c
+++ /dev/null
@@ -1,18 +0,0 @@
-/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
-
-#include <stdio.h>
-#include <openacc.h>
-
-int
-main (int argc, char *argv[])
-{
- int i;
-
-#pragma acc enter data create (i)
- fprintf (stderr, "CheCKpOInT\n");
- acc_create (&i, sizeof i);
-
- return 0;
-}
-
-/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-8.c
libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-8.c
deleted file mode 100644
index e5c0f9cfb323..000000000000
--- libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-8.c
+++ /dev/null
@@ -1,20 +0,0 @@
-/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
-
-#include <stdio.h>
-
-int
-main (int argc, char *argv[])
-{
- int i;
-
-#pragma acc data create (i)
- {
- fprintf (stderr, "CheCKpOInT\n");
-#pragma acc parallel copyin (i)
- ++i;
- }
-
- return 0;
-}
-
-/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
diff --git libgomp/testsuite/libgomp.oacc-fortran/data-already-1.f
libgomp/testsuite/libgomp.oacc-fortran/data-already-1.f
deleted file mode 100644
index fab0ffc99cc3..000000000000
--- libgomp/testsuite/libgomp.oacc-fortran/data-already-1.f
+++ /dev/null
@@ -1,16 +0,0 @@
-! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } }
-
- IMPLICIT NONE
- INCLUDE "openacc_lib.h"
-
- INTEGER I
-
- CALL ACC_COPYIN (I)
- WRITE(0, *) "CheCKpOInT"
-!$ACC DATA COPY (I)
- I = 0
-!$ACC END DATA
-
- END
-
-! { dg-output "CheCKpOInT(\n|\r\n|\r).*" }
diff --git libgomp/testsuite/libgomp.oacc-fortran/data-already-2.f
libgomp/testsuite/libgomp.oacc-fortran/data-already-2.f
deleted file mode 100644
index bd03062670f0..000000000000
--- libgomp/testsuite/libgomp.oacc-fortran/data-already-2.f
+++ /dev/null
@@ -1,16 +0,0 @@
-! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } }
-
- IMPLICIT NONE
-
- INTEGER I
-
-!$ACC DATA PRESENT_OR_COPY (I)
- WRITE(0, *) "CheCKpOInT"
-!$ACC DATA COPYOUT (I)
- I = 0
-!$ACC END DATA
-!$ACC END DATA
-
- END
-
-! { dg-output "CheCKpOInT(\n|\r\n|\r).*" }
diff --git libgomp/testsuite/libgomp.oacc-fortran/data-already-3.f
libgomp/testsuite/libgomp.oacc-fortran/data-already-3.f
deleted file mode 100644
index 60ea3864e4e0..000000000000
--- libgomp/testsuite/libgomp.oacc-fortran/data-already-3.f
+++ /dev/null
@@ -1,15 +0,0 @@
-! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } }
-
- IMPLICIT NONE
- INCLUDE "openacc_lib.h"
-
- INTEGER I
-
-!$ACC DATA PRESENT_OR_COPY (I)
- WRITE(0, *) "CheCKpOInT"
- CALL ACC_COPYIN (I)
-!$ACC END DATA
-
- END
-
-! { dg-output "CheCKpOInT(\n|\r\n|\r).*" }
diff --git libgomp/testsuite/libgomp.oacc-fortran/data-already-4.f
libgomp/testsuite/libgomp.oacc-fortran/data-already-4.f
deleted file mode 100644
index 2abdbf0f868a..000000000000
--- libgomp/testsuite/libgomp.oacc-fortran/data-already-4.f
+++ /dev/null
@@ -1,14 +0,0 @@
-! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } }
-
- IMPLICIT NONE
- INCLUDE "openacc_lib.h"
-
- INTEGER I
-
- CALL ACC_PRESENT_OR_COPYIN (I)
- WRITE(0, *) "CheCKpOInT"
- CALL ACC_COPYIN (I)
-
- END
-
-! { dg-output "CheCKpOInT(\n|\r\n|\r).*" }
diff --git libgomp/testsuite/libgomp.oacc-fortran/data-already-5.f
libgomp/testsuite/libgomp.oacc-fortran/data-already-5.f
deleted file mode 100644
index f361d8c17723..000000000000
--- libgomp/testsuite/libgomp.oacc-fortran/data-already-5.f
+++ /dev/null
@@ -1,14 +0,0 @@
-! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } }
-
- IMPLICIT NONE
- INCLUDE "openacc_lib.h"
-
- INTEGER I
-
-!$ACC ENTER DATA CREATE (I)
- WRITE(0, *) "CheCKpOInT"
- CALL ACC_COPYIN (I)
-
- END
-
-! { dg-output "CheCKpOInT(\n|\r\n|\r).*" }
diff --git libgomp/testsuite/libgomp.oacc-fortran/data-already-6.f
libgomp/testsuite/libgomp.oacc-fortran/data-already-6.f
deleted file mode 100644
index a864737c6923..000000000000
--- libgomp/testsuite/libgomp.oacc-fortran/data-already-6.f
+++ /dev/null
@@ -1,14 +0,0 @@
-! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } }
-
- IMPLICIT NONE
- INCLUDE "openacc_lib.h"
-
- INTEGER I
-
- CALL ACC_PRESENT_OR_COPYIN (I)
- WRITE(0, *) "CheCKpOInT"
-!$ACC ENTER DATA CREATE (I)
-
- END
-
-! { dg-output "CheCKpOInT(\n|\r\n|\r).*" }
diff --git libgomp/testsuite/libgomp.oacc-fortran/data-already-7.f
libgomp/testsuite/libgomp.oacc-fortran/data-already-7.f
deleted file mode 100644
index 0d893280e40b..000000000000
--- libgomp/testsuite/libgomp.oacc-fortran/data-already-7.f
+++ /dev/null
@@ -1,14 +0,0 @@
-! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } }
-
- IMPLICIT NONE
- INCLUDE "openacc_lib.h"
-
- INTEGER I
-
-!$ACC ENTER DATA CREATE (I)
- WRITE(0, *) "CheCKpOInT"
- CALL ACC_CREATE (I)
-
- END
-
-! { dg-output "CheCKpOInT(\n|\r\n|\r).*" }
diff --git libgomp/testsuite/libgomp.oacc-fortran/data-already-8.f
libgomp/testsuite/libgomp.oacc-fortran/data-already-8.f
deleted file mode 100644
index 7a41f67225a4..000000000000
--- libgomp/testsuite/libgomp.oacc-fortran/data-already-8.f
+++ /dev/null
@@ -1,16 +0,0 @@
-! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } }
-
- IMPLICIT NONE
-
- INTEGER I
-
-!$ACC DATA CREATE (I)
- WRITE(0, *) "CheCKpOInT"
-!$ACC PARALLEL COPYIN (I)
- I = 0
-!$ACC END PARALLEL
-!$ACC END DATA
-
- END
-
-! { dg-output "CheCKpOInT(\n|\r\n|\r).*" }
Grüße
Thomas
