diff options
author | Thomas Schwinge <thomas@codesourcery.com> | 2018-12-14 21:43:12 +0100 |
---|---|---|
committer | Thomas Schwinge <tschwinge@gcc.gnu.org> | 2018-12-14 21:43:12 +0100 |
commit | c759830b29fff2288bc0cebfb4b2c479f0b5b30e (patch) | |
tree | 3f358d26336541f0ae0d874adbe278bf6f368943 | |
parent | f847198ec32588beda29a03572a765f9f91b0644 (diff) | |
download | gcc-c759830b29fff2288bc0cebfb4b2c479f0b5b30e.zip gcc-c759830b29fff2288bc0cebfb4b2c479f0b5b30e.tar.gz gcc-c759830b29fff2288bc0cebfb4b2c479f0b5b30e.tar.bz2 |
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.
Co-Authored-By: Chung-Lin Tang <cltang@codesourcery.com>
From-SVN: r267153
19 files changed, 51 insertions, 297 deletions
diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog index 349497d..084f174 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,4 +1,30 @@ 2018-12-14 Thomas Schwinge <thomas@codesourcery.com> + Chung-Lin Tang <cltang@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 <thomas@codesourcery.com> PR libgomp/88495 * plugin/plugin-nvptx.c (nvptx_wait_async): Don't refuse diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c index 72414b7..e352379 100644 --- a/libgomp/oacc-mem.c +++ b/libgomp/oacc-mem.c @@ -544,51 +544,54 @@ acc_create_async (void *h, size_t s, int async) present_create_copy (FLAG_PRESENT | FLAG_CREATE, h, s, async); } -void * -acc_copyin (void *h, size_t s) -{ - return present_create_copy (FLAG_PRESENT | FLAG_CREATE | FLAG_COPY, h, s, - acc_async_sync); -} - -void -acc_copyin_async (void *h, size_t s, int async) -{ - present_create_copy (FLAG_PRESENT | FLAG_CREATE | FLAG_COPY, 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 present_create_copy (FLAG_PRESENT | FLAG_CREATE, h, s, acc_async_sync); + return acc_create (h, s); } -/* 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); + return acc_create (h, s); } #endif void * -acc_present_or_copyin (void *h, size_t s) +acc_copyin (void *h, size_t s) { return present_create_copy (FLAG_PRESENT | FLAG_CREATE | FLAG_COPY, h, s, acc_async_sync); } +void +acc_copyin_async (void *h, size_t s, int async) +{ + present_create_copy (FLAG_PRESENT | FLAG_CREATE | FLAG_COPY, h, s, async); +} + +/* 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 a/libgomp/oacc-parallel.c b/libgomp/oacc-parallel.c index 89b6b6f..e2405d7 100644 --- a/libgomp/oacc-parallel.c +++ b/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 a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-1.c deleted file mode 100644 index fd3b77d..0000000 --- a/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 a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-2.c deleted file mode 100644 index 0118b25..0000000 --- a/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 a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-3.c deleted file mode 100644 index b346c69..0000000 --- a/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 a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-4.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-4.c deleted file mode 100644 index e99ad33..0000000 --- a/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 a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-5.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-5.c deleted file mode 100644 index f8370c0..0000000 --- a/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 a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-6.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-6.c deleted file mode 100644 index d7f4deb..0000000 --- a/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 a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-7.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-7.c deleted file mode 100644 index 54be595..0000000 --- a/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 a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-8.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-8.c deleted file mode 100644 index e5c0f9c..0000000 --- a/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 a/libgomp/testsuite/libgomp.oacc-fortran/data-already-1.f b/libgomp/testsuite/libgomp.oacc-fortran/data-already-1.f deleted file mode 100644 index fab0ffc..0000000 --- a/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 a/libgomp/testsuite/libgomp.oacc-fortran/data-already-2.f b/libgomp/testsuite/libgomp.oacc-fortran/data-already-2.f deleted file mode 100644 index bd03062..0000000 --- a/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 a/libgomp/testsuite/libgomp.oacc-fortran/data-already-3.f b/libgomp/testsuite/libgomp.oacc-fortran/data-already-3.f deleted file mode 100644 index 60ea386..0000000 --- a/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 a/libgomp/testsuite/libgomp.oacc-fortran/data-already-4.f b/libgomp/testsuite/libgomp.oacc-fortran/data-already-4.f deleted file mode 100644 index 2abdbf0..0000000 --- a/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 a/libgomp/testsuite/libgomp.oacc-fortran/data-already-5.f b/libgomp/testsuite/libgomp.oacc-fortran/data-already-5.f deleted file mode 100644 index f361d8c..0000000 --- a/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 a/libgomp/testsuite/libgomp.oacc-fortran/data-already-6.f b/libgomp/testsuite/libgomp.oacc-fortran/data-already-6.f deleted file mode 100644 index a864737..0000000 --- a/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 a/libgomp/testsuite/libgomp.oacc-fortran/data-already-7.f b/libgomp/testsuite/libgomp.oacc-fortran/data-already-7.f deleted file mode 100644 index 0d89328..0000000 --- a/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 a/libgomp/testsuite/libgomp.oacc-fortran/data-already-8.f b/libgomp/testsuite/libgomp.oacc-fortran/data-already-8.f deleted file mode 100644 index 7a41f67..0000000 --- a/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).*" } |