* C/C++ OpenACC: acc_pcopyin, acc_pcreate
@ 2017-05-22 14:28 Thomas Schwinge
2017-05-22 15:23 ` Translate libgomp.oacc-c-c++-common/lib-32.c into Fortran (was: C/C++ OpenACC: acc_pcopyin, acc_pcreate) Thomas Schwinge
2017-05-23 11:08 ` C/C++ OpenACC: acc_pcopyin, acc_pcreate Jakub Jelinek
0 siblings, 2 replies; 8+ messages in thread
From: Thomas Schwinge @ 2017-05-22 14:28 UTC (permalink / raw)
To: gcc-patches, Jakub Jelinek
Hi!
In <openacc.h>, we currently describe acc_pcopyin, acc_pcreate as "old
names", but they're not "old" but really "alternative names", with the
intention to provide them at symbol level, not via "#define"s. OK for
trunk?
commit 89c6599cc343a2f3b087caf9cb47c2c42bb02074
Author: Thomas Schwinge <thomas@codesourcery.com>
Date: Wed Apr 19 15:37:11 2017 +0200
C/C++ OpenACC: acc_pcopyin, acc_pcreate
libgomp/
* openacc.h (acc_pcopyin, acc_pcreate): Provide prototypes instead
of preprocessor definitions.
* libgomp.h (strong_alias): Guard by "#ifdef
HAVE_ATTRIBUTE_ALIAS".
* oacc-mem.c: Provide "acc_pcreate" as alias for
"acc_present_or_create", and "acc_pcopyin" as alias for
"acc_present_or_copyin".
* libgomp.map (OACC_2.0): Add "acc_pcopyin", and "acc_pcreate".
* testsuite/libgomp.oacc-c-c++-common/lib-38.c: Remove, merging
its content into...
* testsuite/libgomp.oacc-c-c++-common/lib-32.c: ... this file.
Extend testing.
---
libgomp/libgomp.h | 5 +-
libgomp/libgomp.map | 2 +
libgomp/oacc-mem.c | 22 ++
libgomp/openacc.h | 7 +-
.../testsuite/libgomp.oacc-c-c++-common/lib-32.c | 241 +++++++++++++++++++--
.../testsuite/libgomp.oacc-c-c++-common/lib-38.c | 64 ------
6 files changed, 254 insertions(+), 87 deletions(-)
diff --git libgomp/libgomp.h libgomp/libgomp.h
index 1769a48..940b5b8 100644
--- libgomp/libgomp.h
+++ libgomp/libgomp.h
@@ -1060,8 +1060,6 @@ extern void gomp_set_nest_lock_25 (omp_nest_lock_25_t *) __GOMP_NOTHROW;
extern void gomp_unset_nest_lock_25 (omp_nest_lock_25_t *) __GOMP_NOTHROW;
extern int gomp_test_nest_lock_25 (omp_nest_lock_25_t *) __GOMP_NOTHROW;
-# define strong_alias(fn, al) \
- extern __typeof (fn) al __attribute__ ((alias (#fn)));
# define omp_lock_symver(fn) \
__asm (".symver g" #fn "_30, " #fn "@@OMP_3.0"); \
__asm (".symver g" #fn "_25, " #fn "@OMP_1.0");
@@ -1085,6 +1083,9 @@ extern int gomp_test_nest_lock_25 (omp_nest_lock_25_t *) __GOMP_NOTHROW;
#endif
#ifdef HAVE_ATTRIBUTE_ALIAS
+# define strong_alias(fn, al) \
+ extern __typeof (fn) al __attribute__ ((alias (#fn)));
+
# define ialias_ulp ialias_str1(__USER_LABEL_PREFIX__)
# define ialias_str1(x) ialias_str2(x)
# define ialias_str2(x) #x
diff --git libgomp/libgomp.map libgomp/libgomp.map
index 4d42c42..c7bf245 100644
--- libgomp/libgomp.map
+++ libgomp/libgomp.map
@@ -335,6 +335,7 @@ OACC_2.0 {
acc_copyin_64_h_;
acc_copyin_array_h_;
acc_present_or_copyin;
+ acc_pcopyin;
acc_present_or_copyin_32_h_;
acc_present_or_copyin_64_h_;
acc_present_or_copyin_array_h_;
@@ -343,6 +344,7 @@ OACC_2.0 {
acc_create_64_h_;
acc_create_array_h_;
acc_present_or_create;
+ acc_pcreate;
acc_present_or_create_32_h_;
acc_present_or_create_64_h_;
acc_present_or_create_array_h_;
diff --git libgomp/oacc-mem.c libgomp/oacc-mem.c
index 2df2202..ff3ed49 100644
--- libgomp/oacc-mem.c
+++ libgomp/oacc-mem.c
@@ -514,12 +514,34 @@ acc_present_or_create (void *h, size_t s)
return present_create_copy (FLAG_PRESENT | FLAG_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);
+}
+#endif
+
void *
acc_present_or_copyin (void *h, size_t s)
{
return present_create_copy (FLAG_PRESENT | FLAG_CREATE | FLAG_COPY, h, s);
}
+/* acc_pcopyin is acc_present_or_copyin by a different name. */
+#ifdef HAVE_ATTRIBUTE_ALIAS
+strong_alias (acc_present_or_copyin, acc_pcopyin)
+#else
+void *
+acc_pcopyin (void *h, size_t s)
+{
+ return acc_present_or_copyin (h, s);
+}
+#endif
+
#define FLAG_COPYOUT (1 << 0)
static void
diff --git libgomp/openacc.h libgomp/openacc.h
index 53d0c39..ebccb18 100644
--- libgomp/openacc.h
+++ libgomp/openacc.h
@@ -91,8 +91,10 @@ void acc_free (void *) __GOACC_NOTHROW;
the standard specifies otherwise. */
void *acc_copyin (void *, size_t) __GOACC_NOTHROW;
void *acc_present_or_copyin (void *, size_t) __GOACC_NOTHROW;
+void *acc_pcopyin (void *, size_t) __GOACC_NOTHROW;
void *acc_create (void *, size_t) __GOACC_NOTHROW;
void *acc_present_or_create (void *, size_t) __GOACC_NOTHROW;
+void *acc_pcreate (void *, size_t) __GOACC_NOTHROW;
void acc_copyout (void *, size_t) __GOACC_NOTHROW;
void acc_delete (void *, size_t) __GOACC_NOTHROW;
void acc_update_device (void *, size_t) __GOACC_NOTHROW;
@@ -105,11 +107,6 @@ int acc_is_present (void *, size_t) __GOACC_NOTHROW;
void acc_memcpy_to_device (void *, void *, size_t) __GOACC_NOTHROW;
void acc_memcpy_from_device (void *, void *, size_t) __GOACC_NOTHROW;
-/* Old names. OpenACC does not specify whether these can or must
- not be macros, inlines or aliases for the new names. */
-#define acc_pcreate acc_present_or_create
-#define acc_pcopyin acc_present_or_copyin
-
/* CUDA-specific routines. */
void *acc_get_current_cuda_device (void) __GOACC_NOTHROW;
void *acc_get_current_cuda_context (void) __GOACC_NOTHROW;
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/lib-32.c libgomp/testsuite/libgomp.oacc-c-c++-common/lib-32.c
index e3f87a8..6a9e995 100644
--- libgomp/testsuite/libgomp.oacc-c-c++-common/lib-32.c
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/lib-32.c
@@ -1,36 +1,245 @@
-/* { dg-do run } */
+/* acc_present_or_create, acc_present_or_copyin, etc. */
+#include <stdbool.h>
#include <stdlib.h>
#include <openacc.h>
int
main (int argc, char **argv)
{
- const int N = 256;
- unsigned char *h;
- void *d1, *d2;
+ int *h, *d;
+ const int N = 10000;
+ const int S = N * sizeof *h;
+ bool shared_mem;
- h = (unsigned char *) malloc (N);
-
- d1 = acc_present_or_create (h, N);
- if (!d1)
+ h = (int *) malloc (S);
+ if (!h)
abort ();
+ for (int i = 0; i < N; ++i)
+ h[i] = i + 0;
- d2 = acc_present_or_create (h, N);
- if (!d2)
- abort ();
+ shared_mem = acc_is_present (h, S);
- if (d1 != d2)
+ d = (int *) acc_present_or_create (h, S);
+ if (!d)
abort ();
+ if (shared_mem)
+ if (h != d)
+ abort ();
+ if (!acc_is_present (h, S))
+ abort ();
+
+#pragma acc parallel loop deviceptr (d)
+ for (int i = 0; i < N; ++i)
+ {
+ d[i] = i + 1;
+ }
+
+ for (int i = 0; i < N; ++i)
+ {
+ if (h[i] != i + (shared_mem ? 1 : 0))
+ abort ();
+ h[i] = i + 2;
+ }
+
+ {
+ int *d_ = (int *) acc_present_or_create (h, S);
+ if (d_ != d)
+ abort ();
+ }
+
+#pragma acc parallel loop deviceptr (d)
+ for (int i = 0; i < N; ++i)
+ {
+ if (d[i] != i + (shared_mem ? 2 : 1))
+ abort ();
+ d[i] = i + 3;
+ }
+
+ for (int i = 0; i < N; ++i)
+ {
+ if (h[i] != i + (shared_mem ? 3 : 2))
+ abort ();
+ h[i] = i + 4;
+ }
+
+ {
+ int *d_ = (int *) acc_pcreate (h, S);
+ if (d_ != d)
+ abort ();
+ }
+
+#pragma acc parallel loop deviceptr (d)
+ for (int i = 0; i < N; ++i)
+ {
+ if (d[i] != i + (shared_mem ? 4 : 3))
+ abort ();
+ d[i] = i + 5;
+ }
+
+ for (int i = 0; i < N; ++i)
+ {
+ if (h[i] != i + (shared_mem ? 5 : 4))
+ abort ();
+ h[i] = i + 6;
+ }
+
+ {
+ int *d_ = (int *) acc_present_or_copyin (h, S);
+ if (d_ != d)
+ abort ();
+ }
+
+#pragma acc parallel loop deviceptr (d)
+ for (int i = 0; i < N; ++i)
+ {
+ if (d[i] != i + (shared_mem ? 6 : 5))
+ abort ();
+ d[i] = i + 7;
+ }
+
+ for (int i = 0; i < N; ++i)
+ {
+ if (h[i] != i + (shared_mem ? 7 : 6))
+ abort ();
+ h[i] = i + 8;
+ }
+
+ {
+ int *d_ = (int *) acc_pcopyin (h, S);
+ if (d_ != d)
+ abort ();
+ }
+
+#pragma acc parallel loop deviceptr (d)
+ for (int i = 0; i < N; ++i)
+ {
+ if (d[i] != i + (shared_mem ? 8 : 7))
+ abort ();
+ d[i] = i + 9;
+ }
- d2 = acc_pcreate (h, N);
- if (!d2)
+ for (int i = 0; i < N; ++i)
+ {
+ if (h[i] != i + (shared_mem ? 9 : 8))
+ abort ();
+ h[i] = i + 10;
+ }
+
+ acc_copyout (h, S);
+ d = NULL;
+ if (!shared_mem)
+ if (acc_is_present (h, S))
+ abort ();
+
+ for (int i = 0; i < N; ++i)
+ {
+ if (h[i] != i + (shared_mem ? 10 : 9))
+ abort ();
+ }
+
+ d = (int *) acc_pcopyin (h, S);
+ if (!d)
+ abort ();
+ if (shared_mem)
+ if (h != d)
+ abort ();
+ if (!acc_is_present (h, S))
abort ();
- if (d1 != d2)
+#pragma acc parallel loop deviceptr (d)
+ for (int i = 0; i < N; ++i)
+ {
+ if (d[i] != i + (shared_mem ? 10 : 9))
+ abort ();
+ d[i] = i + 11;
+ }
+
+ for (int i = 0; i < N; ++i)
+ {
+ if (h[i] != i + (shared_mem ? 11 : 9))
+ abort ();
+ h[i] = i + 12;
+ }
+
+ {
+ int *d_ = (int *) acc_pcopyin (h, S);
+ if (d_ != d)
+ abort ();
+ }
+
+#pragma acc parallel loop deviceptr (d)
+ for (int i = 0; i < N; ++i)
+ {
+ if (d[i] != i + (shared_mem ? 12 : 11))
+ abort ();
+ d[i] = i + 13;
+ }
+
+ for (int i = 0; i < N; ++i)
+ {
+ if (h[i] != i + (shared_mem ? 13 : 12))
+ abort ();
+ h[i] = i + 14;
+ }
+
+ {
+ int *d_ = (int *) acc_pcreate (h, S);
+ if (d_ != d)
+ abort ();
+ }
+
+#pragma acc parallel loop deviceptr (d)
+ for (int i = 0; i < N; ++i)
+ {
+ if (d[i] != i + (shared_mem ? 14 : 13))
+ abort ();
+ d[i] = i + 15;
+ }
+
+ for (int i = 0; i < N; ++i)
+ {
+ if (h[i] != i + (shared_mem ? 15 : 14))
+ abort ();
+ h[i] = i + 16;
+ }
+
+ {
+ int *d_ = (int *) acc_pcreate (h, S);
+ if (d_ != d)
+ abort ();
+ }
+
+#pragma acc parallel loop deviceptr (d)
+ for (int i = 0; i < N; ++i)
+ {
+ if (d[i] != i + (shared_mem ? 16 : 15))
+ abort ();
+ d[i] = i + 17;
+ }
+
+ for (int i = 0; i < N; ++i)
+ {
+ if (h[i] != i + (shared_mem ? 17 : 16))
+ abort ();
+ h[i] = i + 18;
+ }
+
+ acc_update_self (h, S);
+ if (!acc_is_present (h, S))
abort ();
- acc_delete (h, N);
+ for (int i = 0; i < N; ++i)
+ {
+ if (h[i] != i + (shared_mem ? 18 : 17))
+ abort ();
+ }
+
+ acc_delete (h, S);
+ d = NULL;
+ if (!shared_mem)
+ if (acc_is_present (h, S))
+ abort();
free (h);
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/lib-38.c libgomp/testsuite/libgomp.oacc-c-c++-common/lib-38.c
deleted file mode 100644
index 05d8498..0000000
--- libgomp/testsuite/libgomp.oacc-c-c++-common/lib-38.c
+++ /dev/null
@@ -1,64 +0,0 @@
-/* { dg-do run } */
-/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
-
-#include <string.h>
-#include <stdlib.h>
-#include <openacc.h>
-
-int
-main (int argc, char **argv)
-{
- const int N = 256;
- int i;
- unsigned char *h;
- void *d1, *d2;
-
- h = (unsigned char *) malloc (N);
-
- for (i = 0; i < N; i++)
- {
- h[i] = i;
- }
-
- d1 = acc_present_or_copyin (h, N);
- if (!d1)
- abort ();
-
- for (i = 0; i < N; i++)
- {
- h[i] = 0xab;
- }
-
- d2 = acc_present_or_copyin (h, N);
- if (!d2)
- abort ();
-
- if (d1 != d2)
- abort ();
-
- memset (&h[0], 0, N);
-
- acc_copyout (h, N);
-
- for (i = 0; i < N; i++)
- {
- if (h[i] != i)
- abort ();
- }
-
- d2 = acc_pcopyin (h, N);
- if (!d2)
- abort ();
-
- acc_copyout (h, N);
-
- for (i = 0; i < N; i++)
- {
- if (h[i] != i)
- abort ();
- }
-
- free (h);
-
- return 0;
-}
Grüße
Thomas
^ permalink raw reply [flat|nested] 8+ messages in thread
* Translate libgomp.oacc-c-c++-common/lib-32.c into Fortran (was: C/C++ OpenACC: acc_pcopyin, acc_pcreate)
2017-05-22 14:28 C/C++ OpenACC: acc_pcopyin, acc_pcreate Thomas Schwinge
@ 2017-05-22 15:23 ` Thomas Schwinge
2017-05-23 11:12 ` Jakub Jelinek
2017-05-23 11:08 ` C/C++ OpenACC: acc_pcopyin, acc_pcreate Jakub Jelinek
1 sibling, 1 reply; 8+ messages in thread
From: Thomas Schwinge @ 2017-05-22 15:23 UTC (permalink / raw)
To: gcc-patches, Jakub Jelinek
Hi!
On Mon, 22 May 2017 16:26:48 +0200, I wrote:
> C/C++ OpenACC: acc_pcopyin, acc_pcreate
> libgomp/
> [...]
> * testsuite/libgomp.oacc-c-c++-common/lib-38.c: Remove, merging
> its content into...
> * testsuite/libgomp.oacc-c-c++-common/lib-32.c: ... this file.
> Extend testing.
... which I then translated into two Fortran variants; OK for trunk?
commit 04ee44504e6256ee95ae3c6ba91a5960265b3261
Author: Thomas Schwinge <thomas@codesourcery.com>
Date: Wed Apr 19 15:38:51 2017 +0200
Translate libgomp.oacc-c-c++-common/lib-32.c into Fortran
libgomp/
* testsuite/libgomp.oacc-fortran/lib-32-1.f: New file.
* testsuite/libgomp.oacc-fortran/lib-32-2.f: Likewise.
---
.../testsuite/libgomp.oacc-c-c++-common/lib-32.c | 1 +
libgomp/testsuite/libgomp.oacc-fortran/lib-32-1.f | 179 +++++++++++++++++++++
libgomp/testsuite/libgomp.oacc-fortran/lib-32-2.f | 173 ++++++++++++++++++++
3 files changed, 353 insertions(+)
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/lib-32.c libgomp/testsuite/libgomp.oacc-c-c++-common/lib-32.c
index 6a9e995..1696fb6 100644
--- libgomp/testsuite/libgomp.oacc-c-c++-common/lib-32.c
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/lib-32.c
@@ -1,4 +1,5 @@
/* acc_present_or_create, acc_present_or_copyin, etc. */
+/* See also Fortran variants in "../libgomp.oacc-fortran/lib-32*". */
#include <stdbool.h>
#include <stdlib.h>
diff --git libgomp/testsuite/libgomp.oacc-fortran/lib-32-1.f libgomp/testsuite/libgomp.oacc-fortran/lib-32-1.f
new file mode 100644
index 0000000..4606d77
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-fortran/lib-32-1.f
@@ -0,0 +1,179 @@
+! ACC_PRESENT_OR_CREATE, ACC_PRESENT_OR_COPYIN, etc.
+! Variant of "../libgomp.oacc-c-c++-common/lib-32.c".
+! Variant using "openacc_lib.h".
+
+! { dg-do run }
+
+ PROGRAM MAIN
+ IMPLICIT NONE
+ INCLUDE "openacc_lib.h"
+
+ INTEGER, PARAMETER :: N = 10000
+ INTEGER, ALLOCATABLE :: H(:)
+ INTEGER :: I
+ LOGICAL :: SHARED_MEM
+
+ ALLOCATE (H(N))
+ DO I = 1, N
+ H(I) = I + 0
+ END DO
+
+ SHARED_MEM = ACC_IS_PRESENT (H)
+
+ CALL ACC_PRESENT_OR_CREATE (H)
+ IF (.NOT. ACC_IS_PRESENT (H)) CALL ABORT
+
+!$ACC PARALLEL LOOP DEFAULT (PRESENT)
+ DO I = 1, N
+ H(I) = I + 1
+ END DO
+!$ACC END PARALLEL LOOP
+
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (1, 0, SHARED_MEM)) CALL ABORT
+ H(I) = I + 2
+ END DO
+
+ CALL ACC_PRESENT_OR_CREATE (H)
+
+!$ACC PARALLEL LOOP DEFAULT (PRESENT)
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (2, 1, SHARED_MEM)) CALL ABORT
+ H(I) = I + 3
+ END DO
+!$ACC END PARALLEL LOOP
+
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (3, 2, SHARED_MEM)) CALL ABORT
+ H(I) = I + 4
+ END DO
+
+! CALL ACC_PCREATE (H)
+ CALL ACC_PRESENT_OR_CREATE (H)
+
+!$ACC PARALLEL LOOP DEFAULT (PRESENT)
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (4, 3, SHARED_MEM)) CALL ABORT
+ H(I) = I + 5
+ END DO
+!$ACC END PARALLEL LOOP
+
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (5, 4, SHARED_MEM)) CALL ABORT
+ H(I) = I + 6
+ END DO
+
+ CALL ACC_PRESENT_OR_COPYIN (H)
+
+!$ACC PARALLEL LOOP DEFAULT (PRESENT)
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (6, 5, SHARED_MEM)) CALL ABORT
+ H(I) = I + 7
+ END DO
+!$ACC END PARALLEL LOOP
+
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (7, 6, SHARED_MEM)) CALL ABORT
+ H(I) = I + 8
+ END DO
+
+! CALL ACC_PCOPYIN (H)
+ CALL ACC_PRESENT_OR_COPYIN (H)
+
+!$ACC PARALLEL LOOP DEFAULT (PRESENT)
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (8, 7, SHARED_MEM)) CALL ABORT
+ H(I) = I + 9
+ END DO
+!$ACC END PARALLEL LOOP
+
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (9, 8, SHARED_MEM)) CALL ABORT
+ H(I) = I + 10
+ END DO
+
+ CALL ACC_COPYOUT (H)
+ IF (.NOT. SHARED_MEM) THEN
+ IF (ACC_IS_PRESENT (H)) CALL ABORT
+ ENDIF
+
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (10, 9, SHARED_MEM)) CALL ABORT
+ END DO
+
+! CALL ACC_PCOPYIN (H)
+ CALL ACC_PRESENT_OR_COPYIN (H)
+ IF (.NOT. ACC_IS_PRESENT (H)) CALL ABORT
+
+!$ACC PARALLEL LOOP DEFAULT (PRESENT)
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (10, 9, SHARED_MEM)) CALL ABORT
+ H(I) = I + 11
+ END DO
+!$ACC END PARALLEL LOOP
+
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (11, 9, SHARED_MEM)) CALL ABORT
+ H(I) = I + 12
+ END DO
+
+! CALL ACC_PCOPYIN (H)
+ CALL ACC_PRESENT_OR_COPYIN (H)
+
+!$ACC PARALLEL LOOP DEFAULT (PRESENT)
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (12, 11, SHARED_MEM)) CALL ABORT
+ H(I) = I + 13
+ END DO
+!$ACC END PARALLEL LOOP
+
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (13, 12, SHARED_MEM)) CALL ABORT
+ H(I) = I + 14
+ END DO
+
+! CALL ACC_PCREATE (H)
+ CALL ACC_PRESENT_OR_CREATE (H)
+
+!$ACC PARALLEL LOOP DEFAULT (PRESENT)
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (14, 13, SHARED_MEM)) CALL ABORT
+ H(I) = I + 15
+ END DO
+!$ACC END PARALLEL LOOP
+
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (15, 14, SHARED_MEM)) CALL ABORT
+ H(I) = I + 16
+ END DO
+
+! CALL ACC_PCREATE (H)
+ CALL ACC_PRESENT_OR_CREATE (H)
+
+!$ACC PARALLEL LOOP DEFAULT (PRESENT)
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (16, 15, SHARED_MEM)) CALL ABORT
+ H(I) = I + 17
+ END DO
+!$ACC END PARALLEL LOOP
+
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (17, 16, SHARED_MEM)) CALL ABORT
+ H(I) = I + 18
+ END DO
+
+ CALL ACC_UPDATE_SELF (H)
+ IF (.NOT. ACC_IS_PRESENT (H)) CALL ABORT
+
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (18, 17, SHARED_MEM)) CALL ABORT
+ END DO
+
+ CALL ACC_DELETE (H)
+ IF (.NOT. SHARED_MEM) THEN
+ IF (ACC_IS_PRESENT (H)) CALL ABORT
+ ENDIF
+
+ DEALLOCATE (H)
+
+ END PROGRAM MAIN
diff --git libgomp/testsuite/libgomp.oacc-fortran/lib-32-2.f libgomp/testsuite/libgomp.oacc-fortran/lib-32-2.f
new file mode 100644
index 0000000..88f5566
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-fortran/lib-32-2.f
@@ -0,0 +1,173 @@
+! ACC_PRESENT_OR_CREATE, ACC_PRESENT_OR_COPYIN, etc.
+! Variant of "../libgomp.oacc-c-c++-common/lib-32.c".
+! Variant using the "openacc" module.
+
+! { dg-do run }
+
+ PROGRAM MAIN
+ USE OPENACC
+ IMPLICIT NONE
+
+ INTEGER, PARAMETER :: N = 10000
+ INTEGER, ALLOCATABLE :: H(:)
+ INTEGER :: I
+ LOGICAL :: SHARED_MEM
+
+ ALLOCATE (H(N))
+ DO I = 1, N
+ H(I) = I + 0
+ END DO
+
+ SHARED_MEM = ACC_IS_PRESENT (H)
+
+ CALL ACC_PRESENT_OR_CREATE (H)
+ IF (.NOT. ACC_IS_PRESENT (H)) CALL ABORT
+
+!$ACC PARALLEL LOOP DEFAULT (PRESENT)
+ DO I = 1, N
+ H(I) = I + 1
+ END DO
+!$ACC END PARALLEL LOOP
+
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (1, 0, SHARED_MEM)) CALL ABORT
+ H(I) = I + 2
+ END DO
+
+ CALL ACC_PRESENT_OR_CREATE (H)
+
+!$ACC PARALLEL LOOP DEFAULT (PRESENT)
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (2, 1, SHARED_MEM)) CALL ABORT
+ H(I) = I + 3
+ END DO
+!$ACC END PARALLEL LOOP
+
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (3, 2, SHARED_MEM)) CALL ABORT
+ H(I) = I + 4
+ END DO
+
+ CALL ACC_PCREATE (H)
+
+!$ACC PARALLEL LOOP DEFAULT (PRESENT)
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (4, 3, SHARED_MEM)) CALL ABORT
+ H(I) = I + 5
+ END DO
+!$ACC END PARALLEL LOOP
+
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (5, 4, SHARED_MEM)) CALL ABORT
+ H(I) = I + 6
+ END DO
+
+ CALL ACC_PRESENT_OR_COPYIN (H)
+
+!$ACC PARALLEL LOOP DEFAULT (PRESENT)
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (6, 5, SHARED_MEM)) CALL ABORT
+ H(I) = I + 7
+ END DO
+!$ACC END PARALLEL LOOP
+
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (7, 6, SHARED_MEM)) CALL ABORT
+ H(I) = I + 8
+ END DO
+
+ CALL ACC_PCOPYIN (H)
+
+!$ACC PARALLEL LOOP DEFAULT (PRESENT)
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (8, 7, SHARED_MEM)) CALL ABORT
+ H(I) = I + 9
+ END DO
+!$ACC END PARALLEL LOOP
+
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (9, 8, SHARED_MEM)) CALL ABORT
+ H(I) = I + 10
+ END DO
+
+ CALL ACC_COPYOUT (H)
+ IF (.NOT. SHARED_MEM) THEN
+ IF (ACC_IS_PRESENT (H)) CALL ABORT
+ ENDIF
+
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (10, 9, SHARED_MEM)) CALL ABORT
+ END DO
+
+ CALL ACC_PCOPYIN (H)
+ IF (.NOT. ACC_IS_PRESENT (H)) CALL ABORT
+
+!$ACC PARALLEL LOOP DEFAULT (PRESENT)
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (10, 9, SHARED_MEM)) CALL ABORT
+ H(I) = I + 11
+ END DO
+!$ACC END PARALLEL LOOP
+
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (11, 9, SHARED_MEM)) CALL ABORT
+ H(I) = I + 12
+ END DO
+
+ CALL ACC_PCOPYIN (H)
+
+!$ACC PARALLEL LOOP DEFAULT (PRESENT)
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (12, 11, SHARED_MEM)) CALL ABORT
+ H(I) = I + 13
+ END DO
+!$ACC END PARALLEL LOOP
+
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (13, 12, SHARED_MEM)) CALL ABORT
+ H(I) = I + 14
+ END DO
+
+ CALL ACC_PCREATE (H)
+
+!$ACC PARALLEL LOOP DEFAULT (PRESENT)
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (14, 13, SHARED_MEM)) CALL ABORT
+ H(I) = I + 15
+ END DO
+!$ACC END PARALLEL LOOP
+
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (15, 14, SHARED_MEM)) CALL ABORT
+ H(I) = I + 16
+ END DO
+
+ CALL ACC_PCREATE (H)
+
+!$ACC PARALLEL LOOP DEFAULT (PRESENT)
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (16, 15, SHARED_MEM)) CALL ABORT
+ H(I) = I + 17
+ END DO
+!$ACC END PARALLEL LOOP
+
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (17, 16, SHARED_MEM)) CALL ABORT
+ H(I) = I + 18
+ END DO
+
+ CALL ACC_UPDATE_SELF (H)
+ IF (.NOT. ACC_IS_PRESENT (H)) CALL ABORT
+
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (18, 17, SHARED_MEM)) CALL ABORT
+ END DO
+
+ CALL ACC_DELETE (H)
+ IF (.NOT. SHARED_MEM) THEN
+ IF (ACC_IS_PRESENT (H)) CALL ABORT
+ ENDIF
+
+ DEALLOCATE (H)
+
+ END PROGRAM MAIN
Grüße
Thomas
^ permalink raw reply [flat|nested] 8+ messages in thread
* Re: C/C++ OpenACC: acc_pcopyin, acc_pcreate
2017-05-22 14:28 C/C++ OpenACC: acc_pcopyin, acc_pcreate Thomas Schwinge
2017-05-22 15:23 ` Translate libgomp.oacc-c-c++-common/lib-32.c into Fortran (was: C/C++ OpenACC: acc_pcopyin, acc_pcreate) Thomas Schwinge
@ 2017-05-23 11:08 ` Jakub Jelinek
2017-05-24 10:57 ` Thomas Schwinge
1 sibling, 1 reply; 8+ messages in thread
From: Jakub Jelinek @ 2017-05-23 11:08 UTC (permalink / raw)
To: Thomas Schwinge; +Cc: gcc-patches
On Mon, May 22, 2017 at 04:26:48PM +0200, Thomas Schwinge wrote:
> In <openacc.h>, we currently describe acc_pcopyin, acc_pcreate as "old
> names", but they're not "old" but really "alternative names", with the
> intention to provide them at symbol level, not via "#define"s. OK for
> trunk?
No.
> * libgomp.map (OACC_2.0): Add "acc_pcopyin", and "acc_pcreate".
> --- libgomp/libgomp.map
> +++ libgomp/libgomp.map
> @@ -335,6 +335,7 @@ OACC_2.0 {
> acc_copyin_64_h_;
> acc_copyin_array_h_;
> acc_present_or_copyin;
> + acc_pcopyin;
> acc_present_or_copyin_32_h_;
> acc_present_or_copyin_64_h_;
> acc_present_or_copyin_array_h_;
> @@ -343,6 +344,7 @@ OACC_2.0 {
> acc_create_64_h_;
> acc_create_array_h_;
> acc_present_or_create;
> + acc_pcreate;
> acc_present_or_create_32_h_;
> acc_present_or_create_64_h_;
> acc_present_or_create_array_h_;
This is just wrong, new symbols should never be added to an existing symbol
version after a GCC version with that symbol version has been released.
You can add it into OACC_2.0.1, or OACC_1.0, or whatever else, but certainly
not into OACC_2.0.
Another option is just to use something like glibc's sys/cdefs.h
__REDIRECT_NTH macro (including the __USER_LABEL_PREFIX__ stuff)
and just declare those functions as having the asm name of the corresponding
alias. The openacc.h header is for use with GCC only anyway, right?
Unless OpenACC allows one to declare those functions himself and expect
to be able to call them...
Jakub
^ permalink raw reply [flat|nested] 8+ messages in thread
* Re: Translate libgomp.oacc-c-c++-common/lib-32.c into Fortran (was: C/C++ OpenACC: acc_pcopyin, acc_pcreate)
2017-05-22 15:23 ` Translate libgomp.oacc-c-c++-common/lib-32.c into Fortran (was: C/C++ OpenACC: acc_pcopyin, acc_pcreate) Thomas Schwinge
@ 2017-05-23 11:12 ` Jakub Jelinek
2017-05-24 13:45 ` Translate libgomp.oacc-c-c++-common/lib-32.c into Fortran Thomas Schwinge
0 siblings, 1 reply; 8+ messages in thread
From: Jakub Jelinek @ 2017-05-23 11:12 UTC (permalink / raw)
To: Thomas Schwinge; +Cc: gcc-patches
On Mon, May 22, 2017 at 04:29:51PM +0200, Thomas Schwinge wrote:
> Hi!
>
> On Mon, 22 May 2017 16:26:48 +0200, I wrote:
> > C/C++ OpenACC: acc_pcopyin, acc_pcreate
>
> > libgomp/
> > [...]
> > * testsuite/libgomp.oacc-c-c++-common/lib-38.c: Remove, merging
> > its content into...
> > * testsuite/libgomp.oacc-c-c++-common/lib-32.c: ... this file.
> > Extend testing.
>
> ... which I then translated into two Fortran variants; OK for trunk?
>
> commit 04ee44504e6256ee95ae3c6ba91a5960265b3261
> Author: Thomas Schwinge <thomas@codesourcery.com>
> Date: Wed Apr 19 15:38:51 2017 +0200
>
> Translate libgomp.oacc-c-c++-common/lib-32.c into Fortran
>
> libgomp/
> * testsuite/libgomp.oacc-fortran/lib-32-1.f: New file.
> * testsuite/libgomp.oacc-fortran/lib-32-2.f: Likewise.
Ok.
Jakub
^ permalink raw reply [flat|nested] 8+ messages in thread
* Re: C/C++ OpenACC: acc_pcopyin, acc_pcreate
2017-05-23 11:08 ` C/C++ OpenACC: acc_pcopyin, acc_pcreate Jakub Jelinek
@ 2017-05-24 10:57 ` Thomas Schwinge
2017-05-24 11:24 ` Jakub Jelinek
0 siblings, 1 reply; 8+ messages in thread
From: Thomas Schwinge @ 2017-05-24 10:57 UTC (permalink / raw)
To: Jakub Jelinek, gcc-patches
Hi Jakub!
On Tue, 23 May 2017 13:07:15 +0200, Jakub Jelinek <jakub@redhat.com> wrote:
> On Mon, May 22, 2017 at 04:26:48PM +0200, Thomas Schwinge wrote:
> > In <openacc.h>, we currently describe acc_pcopyin, acc_pcreate as "old
> > names", but they're not "old" but really "alternative names", with the
> > intention to provide them at symbol level, not via "#define"s.
> > * libgomp.map (OACC_2.0): Add "acc_pcopyin", and "acc_pcreate".
>
> > --- libgomp/libgomp.map
> > +++ libgomp/libgomp.map
> > @@ -335,6 +335,7 @@ OACC_2.0 {
> > acc_present_or_copyin;
> > + acc_pcopyin;
> > acc_present_or_create;
> > + acc_pcreate;
> This is just wrong, new symbols should never be added to an existing symbol
> version after a GCC version with that symbol version has been released.
> You can add it into OACC_2.0.1, or OACC_1.0, or whatever else, but certainly
> not into OACC_2.0.
I certainly have no problem using a new "OACC_2.0.1" symbol version
instead of "OACC_2.0", but would you please remind me where this
requirement is coming from?
> Another option is just to use something like glibc's sys/cdefs.h
> __REDIRECT_NTH macro (including the __USER_LABEL_PREFIX__ stuff)
> and just declare those functions as having the asm name of the corresponding
> alias. The openacc.h header is for use with GCC only anyway, right?
But that will only redirect them at the user side. The intention here is
to also care for users providing their own declarations instead of using
<openacc.h>, or using "dlsym", and so on -- quite "pathological", I know,
but...
OK for trunk using "OACC_2.0.1" symbol version?
commit 30118ce81354e72e5f32f9ae0ee0f9ef4361747a
Author: Thomas Schwinge <thomas@codesourcery.com>
Date: Wed May 24 12:49:04 2017 +0200
C/C++ OpenACC: acc_pcopyin, acc_pcreate
libgomp/
* openacc.h (acc_pcopyin, acc_pcreate): Provide prototypes instead
of preprocessor definitions.
* libgomp.h (strong_alias): Guard by "#ifdef
HAVE_ATTRIBUTE_ALIAS".
* oacc-mem.c: Provide "acc_pcreate" as alias for
"acc_present_or_create", and "acc_pcopyin" as alias for
"acc_present_or_copyin".
* libgomp.map: New version "OACC_2.0.1".
(OACC_2.0.1): Add "acc_pcopyin", and "acc_pcreate".
* testsuite/libgomp.oacc-c-c++-common/lib-38.c: Remove, merging
its content into...
* testsuite/libgomp.oacc-c-c++-common/lib-32.c: ... this file.
Extend testing.
---
libgomp/libgomp.h | 5 +-
libgomp/libgomp.map | 6 +
libgomp/oacc-mem.c | 22 ++
libgomp/openacc.h | 7 +-
.../testsuite/libgomp.oacc-c-c++-common/lib-32.c | 241 +++++++++++++++++++--
.../testsuite/libgomp.oacc-c-c++-common/lib-38.c | 64 ------
6 files changed, 258 insertions(+), 87 deletions(-)
diff --git libgomp/libgomp.h libgomp/libgomp.h
index 1769a48..940b5b8 100644
--- libgomp/libgomp.h
+++ libgomp/libgomp.h
@@ -1060,8 +1060,6 @@ extern void gomp_set_nest_lock_25 (omp_nest_lock_25_t *) __GOMP_NOTHROW;
extern void gomp_unset_nest_lock_25 (omp_nest_lock_25_t *) __GOMP_NOTHROW;
extern int gomp_test_nest_lock_25 (omp_nest_lock_25_t *) __GOMP_NOTHROW;
-# define strong_alias(fn, al) \
- extern __typeof (fn) al __attribute__ ((alias (#fn)));
# define omp_lock_symver(fn) \
__asm (".symver g" #fn "_30, " #fn "@@OMP_3.0"); \
__asm (".symver g" #fn "_25, " #fn "@OMP_1.0");
@@ -1085,6 +1083,9 @@ extern int gomp_test_nest_lock_25 (omp_nest_lock_25_t *) __GOMP_NOTHROW;
#endif
#ifdef HAVE_ATTRIBUTE_ALIAS
+# define strong_alias(fn, al) \
+ extern __typeof (fn) al __attribute__ ((alias (#fn)));
+
# define ialias_ulp ialias_str1(__USER_LABEL_PREFIX__)
# define ialias_str1(x) ialias_str2(x)
# define ialias_str2(x) #x
diff --git libgomp/libgomp.map libgomp/libgomp.map
index 4d42c42..b43c6de 100644
--- libgomp/libgomp.map
+++ libgomp/libgomp.map
@@ -378,6 +378,12 @@ OACC_2.0 {
acc_set_cuda_stream;
};
+OACC_2.0.1 {
+ global:
+ acc_pcopyin;
+ acc_pcreate;
+} OACC_2.0;
+
GOACC_2.0 {
global:
GOACC_data_end;
diff --git libgomp/oacc-mem.c libgomp/oacc-mem.c
index 2df2202..ff3ed49 100644
--- libgomp/oacc-mem.c
+++ libgomp/oacc-mem.c
@@ -514,12 +514,34 @@ acc_present_or_create (void *h, size_t s)
return present_create_copy (FLAG_PRESENT | FLAG_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);
+}
+#endif
+
void *
acc_present_or_copyin (void *h, size_t s)
{
return present_create_copy (FLAG_PRESENT | FLAG_CREATE | FLAG_COPY, h, s);
}
+/* acc_pcopyin is acc_present_or_copyin by a different name. */
+#ifdef HAVE_ATTRIBUTE_ALIAS
+strong_alias (acc_present_or_copyin, acc_pcopyin)
+#else
+void *
+acc_pcopyin (void *h, size_t s)
+{
+ return acc_present_or_copyin (h, s);
+}
+#endif
+
#define FLAG_COPYOUT (1 << 0)
static void
diff --git libgomp/openacc.h libgomp/openacc.h
index 53d0c39..ebccb18 100644
--- libgomp/openacc.h
+++ libgomp/openacc.h
@@ -91,8 +91,10 @@ void acc_free (void *) __GOACC_NOTHROW;
the standard specifies otherwise. */
void *acc_copyin (void *, size_t) __GOACC_NOTHROW;
void *acc_present_or_copyin (void *, size_t) __GOACC_NOTHROW;
+void *acc_pcopyin (void *, size_t) __GOACC_NOTHROW;
void *acc_create (void *, size_t) __GOACC_NOTHROW;
void *acc_present_or_create (void *, size_t) __GOACC_NOTHROW;
+void *acc_pcreate (void *, size_t) __GOACC_NOTHROW;
void acc_copyout (void *, size_t) __GOACC_NOTHROW;
void acc_delete (void *, size_t) __GOACC_NOTHROW;
void acc_update_device (void *, size_t) __GOACC_NOTHROW;
@@ -105,11 +107,6 @@ int acc_is_present (void *, size_t) __GOACC_NOTHROW;
void acc_memcpy_to_device (void *, void *, size_t) __GOACC_NOTHROW;
void acc_memcpy_from_device (void *, void *, size_t) __GOACC_NOTHROW;
-/* Old names. OpenACC does not specify whether these can or must
- not be macros, inlines or aliases for the new names. */
-#define acc_pcreate acc_present_or_create
-#define acc_pcopyin acc_present_or_copyin
-
/* CUDA-specific routines. */
void *acc_get_current_cuda_device (void) __GOACC_NOTHROW;
void *acc_get_current_cuda_context (void) __GOACC_NOTHROW;
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/lib-32.c libgomp/testsuite/libgomp.oacc-c-c++-common/lib-32.c
index e3f87a8..6a9e995 100644
--- libgomp/testsuite/libgomp.oacc-c-c++-common/lib-32.c
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/lib-32.c
@@ -1,36 +1,245 @@
-/* { dg-do run } */
+/* acc_present_or_create, acc_present_or_copyin, etc. */
+#include <stdbool.h>
#include <stdlib.h>
#include <openacc.h>
int
main (int argc, char **argv)
{
- const int N = 256;
- unsigned char *h;
- void *d1, *d2;
+ int *h, *d;
+ const int N = 10000;
+ const int S = N * sizeof *h;
+ bool shared_mem;
- h = (unsigned char *) malloc (N);
-
- d1 = acc_present_or_create (h, N);
- if (!d1)
+ h = (int *) malloc (S);
+ if (!h)
abort ();
+ for (int i = 0; i < N; ++i)
+ h[i] = i + 0;
- d2 = acc_present_or_create (h, N);
- if (!d2)
- abort ();
+ shared_mem = acc_is_present (h, S);
- if (d1 != d2)
+ d = (int *) acc_present_or_create (h, S);
+ if (!d)
abort ();
+ if (shared_mem)
+ if (h != d)
+ abort ();
+ if (!acc_is_present (h, S))
+ abort ();
+
+#pragma acc parallel loop deviceptr (d)
+ for (int i = 0; i < N; ++i)
+ {
+ d[i] = i + 1;
+ }
+
+ for (int i = 0; i < N; ++i)
+ {
+ if (h[i] != i + (shared_mem ? 1 : 0))
+ abort ();
+ h[i] = i + 2;
+ }
+
+ {
+ int *d_ = (int *) acc_present_or_create (h, S);
+ if (d_ != d)
+ abort ();
+ }
+
+#pragma acc parallel loop deviceptr (d)
+ for (int i = 0; i < N; ++i)
+ {
+ if (d[i] != i + (shared_mem ? 2 : 1))
+ abort ();
+ d[i] = i + 3;
+ }
+
+ for (int i = 0; i < N; ++i)
+ {
+ if (h[i] != i + (shared_mem ? 3 : 2))
+ abort ();
+ h[i] = i + 4;
+ }
+
+ {
+ int *d_ = (int *) acc_pcreate (h, S);
+ if (d_ != d)
+ abort ();
+ }
+
+#pragma acc parallel loop deviceptr (d)
+ for (int i = 0; i < N; ++i)
+ {
+ if (d[i] != i + (shared_mem ? 4 : 3))
+ abort ();
+ d[i] = i + 5;
+ }
+
+ for (int i = 0; i < N; ++i)
+ {
+ if (h[i] != i + (shared_mem ? 5 : 4))
+ abort ();
+ h[i] = i + 6;
+ }
+
+ {
+ int *d_ = (int *) acc_present_or_copyin (h, S);
+ if (d_ != d)
+ abort ();
+ }
+
+#pragma acc parallel loop deviceptr (d)
+ for (int i = 0; i < N; ++i)
+ {
+ if (d[i] != i + (shared_mem ? 6 : 5))
+ abort ();
+ d[i] = i + 7;
+ }
+
+ for (int i = 0; i < N; ++i)
+ {
+ if (h[i] != i + (shared_mem ? 7 : 6))
+ abort ();
+ h[i] = i + 8;
+ }
+
+ {
+ int *d_ = (int *) acc_pcopyin (h, S);
+ if (d_ != d)
+ abort ();
+ }
+
+#pragma acc parallel loop deviceptr (d)
+ for (int i = 0; i < N; ++i)
+ {
+ if (d[i] != i + (shared_mem ? 8 : 7))
+ abort ();
+ d[i] = i + 9;
+ }
- d2 = acc_pcreate (h, N);
- if (!d2)
+ for (int i = 0; i < N; ++i)
+ {
+ if (h[i] != i + (shared_mem ? 9 : 8))
+ abort ();
+ h[i] = i + 10;
+ }
+
+ acc_copyout (h, S);
+ d = NULL;
+ if (!shared_mem)
+ if (acc_is_present (h, S))
+ abort ();
+
+ for (int i = 0; i < N; ++i)
+ {
+ if (h[i] != i + (shared_mem ? 10 : 9))
+ abort ();
+ }
+
+ d = (int *) acc_pcopyin (h, S);
+ if (!d)
+ abort ();
+ if (shared_mem)
+ if (h != d)
+ abort ();
+ if (!acc_is_present (h, S))
abort ();
- if (d1 != d2)
+#pragma acc parallel loop deviceptr (d)
+ for (int i = 0; i < N; ++i)
+ {
+ if (d[i] != i + (shared_mem ? 10 : 9))
+ abort ();
+ d[i] = i + 11;
+ }
+
+ for (int i = 0; i < N; ++i)
+ {
+ if (h[i] != i + (shared_mem ? 11 : 9))
+ abort ();
+ h[i] = i + 12;
+ }
+
+ {
+ int *d_ = (int *) acc_pcopyin (h, S);
+ if (d_ != d)
+ abort ();
+ }
+
+#pragma acc parallel loop deviceptr (d)
+ for (int i = 0; i < N; ++i)
+ {
+ if (d[i] != i + (shared_mem ? 12 : 11))
+ abort ();
+ d[i] = i + 13;
+ }
+
+ for (int i = 0; i < N; ++i)
+ {
+ if (h[i] != i + (shared_mem ? 13 : 12))
+ abort ();
+ h[i] = i + 14;
+ }
+
+ {
+ int *d_ = (int *) acc_pcreate (h, S);
+ if (d_ != d)
+ abort ();
+ }
+
+#pragma acc parallel loop deviceptr (d)
+ for (int i = 0; i < N; ++i)
+ {
+ if (d[i] != i + (shared_mem ? 14 : 13))
+ abort ();
+ d[i] = i + 15;
+ }
+
+ for (int i = 0; i < N; ++i)
+ {
+ if (h[i] != i + (shared_mem ? 15 : 14))
+ abort ();
+ h[i] = i + 16;
+ }
+
+ {
+ int *d_ = (int *) acc_pcreate (h, S);
+ if (d_ != d)
+ abort ();
+ }
+
+#pragma acc parallel loop deviceptr (d)
+ for (int i = 0; i < N; ++i)
+ {
+ if (d[i] != i + (shared_mem ? 16 : 15))
+ abort ();
+ d[i] = i + 17;
+ }
+
+ for (int i = 0; i < N; ++i)
+ {
+ if (h[i] != i + (shared_mem ? 17 : 16))
+ abort ();
+ h[i] = i + 18;
+ }
+
+ acc_update_self (h, S);
+ if (!acc_is_present (h, S))
abort ();
- acc_delete (h, N);
+ for (int i = 0; i < N; ++i)
+ {
+ if (h[i] != i + (shared_mem ? 18 : 17))
+ abort ();
+ }
+
+ acc_delete (h, S);
+ d = NULL;
+ if (!shared_mem)
+ if (acc_is_present (h, S))
+ abort();
free (h);
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/lib-38.c libgomp/testsuite/libgomp.oacc-c-c++-common/lib-38.c
deleted file mode 100644
index 05d8498..0000000
--- libgomp/testsuite/libgomp.oacc-c-c++-common/lib-38.c
+++ /dev/null
@@ -1,64 +0,0 @@
-/* { dg-do run } */
-/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
-
-#include <string.h>
-#include <stdlib.h>
-#include <openacc.h>
-
-int
-main (int argc, char **argv)
-{
- const int N = 256;
- int i;
- unsigned char *h;
- void *d1, *d2;
-
- h = (unsigned char *) malloc (N);
-
- for (i = 0; i < N; i++)
- {
- h[i] = i;
- }
-
- d1 = acc_present_or_copyin (h, N);
- if (!d1)
- abort ();
-
- for (i = 0; i < N; i++)
- {
- h[i] = 0xab;
- }
-
- d2 = acc_present_or_copyin (h, N);
- if (!d2)
- abort ();
-
- if (d1 != d2)
- abort ();
-
- memset (&h[0], 0, N);
-
- acc_copyout (h, N);
-
- for (i = 0; i < N; i++)
- {
- if (h[i] != i)
- abort ();
- }
-
- d2 = acc_pcopyin (h, N);
- if (!d2)
- abort ();
-
- acc_copyout (h, N);
-
- for (i = 0; i < N; i++)
- {
- if (h[i] != i)
- abort ();
- }
-
- free (h);
-
- return 0;
-}
Grüße
Thomas
^ permalink raw reply [flat|nested] 8+ messages in thread
* Re: C/C++ OpenACC: acc_pcopyin, acc_pcreate
2017-05-24 10:57 ` Thomas Schwinge
@ 2017-05-24 11:24 ` Jakub Jelinek
2017-05-24 13:41 ` Thomas Schwinge
0 siblings, 1 reply; 8+ messages in thread
From: Jakub Jelinek @ 2017-05-24 11:24 UTC (permalink / raw)
To: Thomas Schwinge; +Cc: gcc-patches
On Wed, May 24, 2017 at 12:55:27PM +0200, Thomas Schwinge wrote:
> Hi Jakub!
>
> On Tue, 23 May 2017 13:07:15 +0200, Jakub Jelinek <jakub@redhat.com> wrote:
> > On Mon, May 22, 2017 at 04:26:48PM +0200, Thomas Schwinge wrote:
> > > In <openacc.h>, we currently describe acc_pcopyin, acc_pcreate as "old
> > > names", but they're not "old" but really "alternative names", with the
> > > intention to provide them at symbol level, not via "#define"s.
>
> > > * libgomp.map (OACC_2.0): Add "acc_pcopyin", and "acc_pcreate".
> >
> > > --- libgomp/libgomp.map
> > > +++ libgomp/libgomp.map
> > > @@ -335,6 +335,7 @@ OACC_2.0 {
>
> > > acc_present_or_copyin;
> > > + acc_pcopyin;
>
> > > acc_present_or_create;
> > > + acc_pcreate;
>
> > This is just wrong, new symbols should never be added to an existing symbol
> > version after a GCC version with that symbol version has been released.
> > You can add it into OACC_2.0.1, or OACC_1.0, or whatever else, but certainly
> > not into OACC_2.0.
>
> I certainly have no problem using a new "OACC_2.0.1" symbol version
> instead of "OACC_2.0", but would you please remind me where this
> requirement is coming from?
There are many reasons, one is e.g. that the dynamic linker uses the symbol
version names during library loading to verify the library satisfies the
needs of a binary or some other library; if you add further symbols into
existing symbol version, then it will be satisfied even by library that
doesn't provide the symbol (in the above case say acc_pcopyin), and you can
get a failure either much later (during relocation processing, with lazy
binding that can be much later), or not at all (with lazy binding and if
the function isn't called at all during the runtime of the program).
Another one is e.g. rpm uses symbol versions in dependencies, but only
those and not at the required/provided symbol granularity. So, if you
have acc_pcopyin@@OACC_2.0.1 and use it in some program, the program
will have a dependency on libgomp.so.1()(OACC_2.0.1) and thus you
automatically get the right library that provides it. If you add to an
existing symbol version, the dependency will be satisfied even for libraries
that don't provide it at all.
All in all, symbol versions aren't there just for documentation purposes,
but serve various other roles, including the possibility to change ABI
of existing symbols, but not limited to that. The policy of not adding
new symbols into symbol versions that were already shipped in released
libraries is something that is strictly followed in glibc, gcc libraries
(I remember in libstdc++ a bug happened some years ago and some symbols
were unintentionally exported at an earlier version, which has been cured
by moving into a new symbol version and providing the old one just as an @
(vs. @@) symver for compatibility only; e.g. the current libstdc++ check-abi
stuff verifies this stuff).
For more info see e.g. https://www.akkadia.org/drepper/symbol-versioning
> But that will only redirect them at the user side. The intention here is
The question is if the OpenACC standard allows that or not.
> to also care for users providing their own declarations instead of using
> <openacc.h>, or using "dlsym", and so on -- quite "pathological", I know,
> but...
>
> OK for trunk using "OACC_2.0.1" symbol version?
Ok.
> commit 30118ce81354e72e5f32f9ae0ee0f9ef4361747a
> Author: Thomas Schwinge <thomas@codesourcery.com>
> Date: Wed May 24 12:49:04 2017 +0200
>
> C/C++ OpenACC: acc_pcopyin, acc_pcreate
>
> libgomp/
> * openacc.h (acc_pcopyin, acc_pcreate): Provide prototypes instead
> of preprocessor definitions.
> * libgomp.h (strong_alias): Guard by "#ifdef
> HAVE_ATTRIBUTE_ALIAS".
> * oacc-mem.c: Provide "acc_pcreate" as alias for
> "acc_present_or_create", and "acc_pcopyin" as alias for
> "acc_present_or_copyin".
> * libgomp.map: New version "OACC_2.0.1".
> (OACC_2.0.1): Add "acc_pcopyin", and "acc_pcreate".
> * testsuite/libgomp.oacc-c-c++-common/lib-38.c: Remove, merging
> its content into...
> * testsuite/libgomp.oacc-c-c++-common/lib-32.c: ... this file.
> Extend testing.
Jakub
^ permalink raw reply [flat|nested] 8+ messages in thread
* Re: C/C++ OpenACC: acc_pcopyin, acc_pcreate
2017-05-24 11:24 ` Jakub Jelinek
@ 2017-05-24 13:41 ` Thomas Schwinge
0 siblings, 0 replies; 8+ messages in thread
From: Thomas Schwinge @ 2017-05-24 13:41 UTC (permalink / raw)
To: Jakub Jelinek, gcc-patches
Hi!
On Wed, 24 May 2017 13:20:54 +0200, Jakub Jelinek <jakub@redhat.com> wrote:
> On Wed, May 24, 2017 at 12:55:27PM +0200, Thomas Schwinge wrote:
> > On Tue, 23 May 2017 13:07:15 +0200, Jakub Jelinek <jakub@redhat.com> wrote:
> > > new symbols should never be added to an existing symbol
> > > version after a GCC version with that symbol version has been released.
> > would you please remind me where this
> > requirement is coming from?
>
> There are many reasons
Thanks for the explanation!
> one is e.g. that the dynamic linker uses the symbol
> version names during library loading to verify the library satisfies the
> needs of a binary or some other library; if you add further symbols into
> existing symbol version, then it will be satisfied even by library that
> doesn't provide the symbol (in the above case say acc_pcopyin), and you can
> get a failure either much later (during relocation processing, with lazy
> binding that can be much later), or not at all (with lazy binding and if
> the function isn't called at all during the runtime of the program).
Some might consider that a feature. ;-)
> Another one is e.g. rpm uses symbol versions in dependencies, but only
> those and not at the required/provided symbol granularity. So, if you
> have acc_pcopyin@@OACC_2.0.1 and use it in some program, the program
> will have a dependency on libgomp.so.1()(OACC_2.0.1) and thus you
> automatically get the right library that provides it. If you add to an
> existing symbol version, the dependency will be satisfied even for libraries
> that don't provide it at all.
Ah, yes, I see -- a fine reason.
> All in all, symbol versions aren't there just for documentation purposes,
> but serve various other roles, including the possibility to change ABI
> of existing symbols, but not limited to that.
Sure, I was aware of that specific use -- but that's not relevant to the
case discussed here: a symbol's (non-)existence will always be detected
reliably, whether lazily (see above), or at load time. That's where I
had gotten the idea from that it's fine to use existing versions for
these new symbols. (But you reminded me why it's better to use a new
version here, too.)
> The policy of not adding
> new symbols into symbol versions that were already shipped in released
> libraries is something that is strictly followed in glibc, gcc libraries
> (I remember in libstdc++ a bug happened some years ago and some symbols
> were unintentionally exported at an earlier version, which has been cured
> by moving into a new symbol version and providing the old one just as an @
> (vs. @@) symver for compatibility only; e.g. the current libstdc++ check-abi
> stuff verifies this stuff).
> For more info see e.g. https://www.akkadia.org/drepper/symbol-versioning
Yeah, thanks, I had read that a few years ago.
> > But that will only redirect them at the user side. The intention here is
>
> The question is if the OpenACC standard allows that or not.
Let's put it this way: it doesn't explicitly disallow that.
> > to also care for users providing their own declarations instead of using
> > <openacc.h>, or using "dlsym", and so on -- quite "pathological", I know,
> > but...
> > OK for trunk using "OACC_2.0.1" symbol version?
>
> Ok.
Thanks. As posted, committed to trunk in r248410:
commit e4d15e02dfbaf61fb5655595423baaff4a345a60
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Wed May 24 13:23:34 2017 +0000
C/C++ OpenACC: acc_pcopyin, acc_pcreate
libgomp/
* openacc.h (acc_pcopyin, acc_pcreate): Provide prototypes instead
of preprocessor definitions.
* libgomp.h (strong_alias): Guard by "#ifdef
HAVE_ATTRIBUTE_ALIAS".
* oacc-mem.c: Provide "acc_pcreate" as alias for
"acc_present_or_create", and "acc_pcopyin" as alias for
"acc_present_or_copyin".
* libgomp.map: New version "OACC_2.0.1".
(OACC_2.0.1): Add "acc_pcopyin", and "acc_pcreate".
* testsuite/libgomp.oacc-c-c++-common/lib-38.c: Remove, merging
its content into...
* testsuite/libgomp.oacc-c-c++-common/lib-32.c: ... this file.
Extend testing.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@248410 138bc75d-0d04-0410-961f-82ee72b054a4
---
libgomp/ChangeLog | 14 ++
libgomp/libgomp.h | 5 +-
libgomp/libgomp.map | 6 +
libgomp/oacc-mem.c | 22 ++
libgomp/openacc.h | 7 +-
.../testsuite/libgomp.oacc-c-c++-common/lib-32.c | 241 +++++++++++++++++++--
.../testsuite/libgomp.oacc-c-c++-common/lib-38.c | 64 ------
7 files changed, 272 insertions(+), 87 deletions(-)
diff --git libgomp/ChangeLog libgomp/ChangeLog
index 14e95ef..4a95755 100644
--- libgomp/ChangeLog
+++ libgomp/ChangeLog
@@ -1,5 +1,19 @@
2017-05-24 Thomas Schwinge <thomas@codesourcery.com>
+ * openacc.h (acc_pcopyin, acc_pcreate): Provide prototypes instead
+ of preprocessor definitions.
+ * libgomp.h (strong_alias): Guard by "#ifdef
+ HAVE_ATTRIBUTE_ALIAS".
+ * oacc-mem.c: Provide "acc_pcreate" as alias for
+ "acc_present_or_create", and "acc_pcopyin" as alias for
+ "acc_present_or_copyin".
+ * libgomp.map: New version "OACC_2.0.1".
+ (OACC_2.0.1): Add "acc_pcopyin", and "acc_pcreate".
+ * testsuite/libgomp.oacc-c-c++-common/lib-38.c: Remove, merging
+ its content into...
+ * testsuite/libgomp.oacc-c-c++-common/lib-32.c: ... this file.
+ Extend testing.
+
* plugin/plugin-nvptx.c (nvptx_get_num_devices): Debugging output
when disabling nvptx offloading.
diff --git libgomp/libgomp.h libgomp/libgomp.h
index 1769a48..940b5b8 100644
--- libgomp/libgomp.h
+++ libgomp/libgomp.h
@@ -1060,8 +1060,6 @@ extern void gomp_set_nest_lock_25 (omp_nest_lock_25_t *) __GOMP_NOTHROW;
extern void gomp_unset_nest_lock_25 (omp_nest_lock_25_t *) __GOMP_NOTHROW;
extern int gomp_test_nest_lock_25 (omp_nest_lock_25_t *) __GOMP_NOTHROW;
-# define strong_alias(fn, al) \
- extern __typeof (fn) al __attribute__ ((alias (#fn)));
# define omp_lock_symver(fn) \
__asm (".symver g" #fn "_30, " #fn "@@OMP_3.0"); \
__asm (".symver g" #fn "_25, " #fn "@OMP_1.0");
@@ -1085,6 +1083,9 @@ extern int gomp_test_nest_lock_25 (omp_nest_lock_25_t *) __GOMP_NOTHROW;
#endif
#ifdef HAVE_ATTRIBUTE_ALIAS
+# define strong_alias(fn, al) \
+ extern __typeof (fn) al __attribute__ ((alias (#fn)));
+
# define ialias_ulp ialias_str1(__USER_LABEL_PREFIX__)
# define ialias_str1(x) ialias_str2(x)
# define ialias_str2(x) #x
diff --git libgomp/libgomp.map libgomp/libgomp.map
index 4d42c42..b43c6de 100644
--- libgomp/libgomp.map
+++ libgomp/libgomp.map
@@ -378,6 +378,12 @@ OACC_2.0 {
acc_set_cuda_stream;
};
+OACC_2.0.1 {
+ global:
+ acc_pcopyin;
+ acc_pcreate;
+} OACC_2.0;
+
GOACC_2.0 {
global:
GOACC_data_end;
diff --git libgomp/oacc-mem.c libgomp/oacc-mem.c
index 2df2202..ff3ed49 100644
--- libgomp/oacc-mem.c
+++ libgomp/oacc-mem.c
@@ -514,12 +514,34 @@ acc_present_or_create (void *h, size_t s)
return present_create_copy (FLAG_PRESENT | FLAG_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);
+}
+#endif
+
void *
acc_present_or_copyin (void *h, size_t s)
{
return present_create_copy (FLAG_PRESENT | FLAG_CREATE | FLAG_COPY, h, s);
}
+/* acc_pcopyin is acc_present_or_copyin by a different name. */
+#ifdef HAVE_ATTRIBUTE_ALIAS
+strong_alias (acc_present_or_copyin, acc_pcopyin)
+#else
+void *
+acc_pcopyin (void *h, size_t s)
+{
+ return acc_present_or_copyin (h, s);
+}
+#endif
+
#define FLAG_COPYOUT (1 << 0)
static void
diff --git libgomp/openacc.h libgomp/openacc.h
index 53d0c39..ebccb18 100644
--- libgomp/openacc.h
+++ libgomp/openacc.h
@@ -91,8 +91,10 @@ void acc_free (void *) __GOACC_NOTHROW;
the standard specifies otherwise. */
void *acc_copyin (void *, size_t) __GOACC_NOTHROW;
void *acc_present_or_copyin (void *, size_t) __GOACC_NOTHROW;
+void *acc_pcopyin (void *, size_t) __GOACC_NOTHROW;
void *acc_create (void *, size_t) __GOACC_NOTHROW;
void *acc_present_or_create (void *, size_t) __GOACC_NOTHROW;
+void *acc_pcreate (void *, size_t) __GOACC_NOTHROW;
void acc_copyout (void *, size_t) __GOACC_NOTHROW;
void acc_delete (void *, size_t) __GOACC_NOTHROW;
void acc_update_device (void *, size_t) __GOACC_NOTHROW;
@@ -105,11 +107,6 @@ int acc_is_present (void *, size_t) __GOACC_NOTHROW;
void acc_memcpy_to_device (void *, void *, size_t) __GOACC_NOTHROW;
void acc_memcpy_from_device (void *, void *, size_t) __GOACC_NOTHROW;
-/* Old names. OpenACC does not specify whether these can or must
- not be macros, inlines or aliases for the new names. */
-#define acc_pcreate acc_present_or_create
-#define acc_pcopyin acc_present_or_copyin
-
/* CUDA-specific routines. */
void *acc_get_current_cuda_device (void) __GOACC_NOTHROW;
void *acc_get_current_cuda_context (void) __GOACC_NOTHROW;
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/lib-32.c libgomp/testsuite/libgomp.oacc-c-c++-common/lib-32.c
index e3f87a8..6a9e995 100644
--- libgomp/testsuite/libgomp.oacc-c-c++-common/lib-32.c
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/lib-32.c
@@ -1,36 +1,245 @@
-/* { dg-do run } */
+/* acc_present_or_create, acc_present_or_copyin, etc. */
+#include <stdbool.h>
#include <stdlib.h>
#include <openacc.h>
int
main (int argc, char **argv)
{
- const int N = 256;
- unsigned char *h;
- void *d1, *d2;
+ int *h, *d;
+ const int N = 10000;
+ const int S = N * sizeof *h;
+ bool shared_mem;
- h = (unsigned char *) malloc (N);
-
- d1 = acc_present_or_create (h, N);
- if (!d1)
+ h = (int *) malloc (S);
+ if (!h)
abort ();
+ for (int i = 0; i < N; ++i)
+ h[i] = i + 0;
- d2 = acc_present_or_create (h, N);
- if (!d2)
- abort ();
+ shared_mem = acc_is_present (h, S);
- if (d1 != d2)
+ d = (int *) acc_present_or_create (h, S);
+ if (!d)
abort ();
+ if (shared_mem)
+ if (h != d)
+ abort ();
+ if (!acc_is_present (h, S))
+ abort ();
+
+#pragma acc parallel loop deviceptr (d)
+ for (int i = 0; i < N; ++i)
+ {
+ d[i] = i + 1;
+ }
+
+ for (int i = 0; i < N; ++i)
+ {
+ if (h[i] != i + (shared_mem ? 1 : 0))
+ abort ();
+ h[i] = i + 2;
+ }
+
+ {
+ int *d_ = (int *) acc_present_or_create (h, S);
+ if (d_ != d)
+ abort ();
+ }
+
+#pragma acc parallel loop deviceptr (d)
+ for (int i = 0; i < N; ++i)
+ {
+ if (d[i] != i + (shared_mem ? 2 : 1))
+ abort ();
+ d[i] = i + 3;
+ }
+
+ for (int i = 0; i < N; ++i)
+ {
+ if (h[i] != i + (shared_mem ? 3 : 2))
+ abort ();
+ h[i] = i + 4;
+ }
+
+ {
+ int *d_ = (int *) acc_pcreate (h, S);
+ if (d_ != d)
+ abort ();
+ }
+
+#pragma acc parallel loop deviceptr (d)
+ for (int i = 0; i < N; ++i)
+ {
+ if (d[i] != i + (shared_mem ? 4 : 3))
+ abort ();
+ d[i] = i + 5;
+ }
+
+ for (int i = 0; i < N; ++i)
+ {
+ if (h[i] != i + (shared_mem ? 5 : 4))
+ abort ();
+ h[i] = i + 6;
+ }
+
+ {
+ int *d_ = (int *) acc_present_or_copyin (h, S);
+ if (d_ != d)
+ abort ();
+ }
+
+#pragma acc parallel loop deviceptr (d)
+ for (int i = 0; i < N; ++i)
+ {
+ if (d[i] != i + (shared_mem ? 6 : 5))
+ abort ();
+ d[i] = i + 7;
+ }
+
+ for (int i = 0; i < N; ++i)
+ {
+ if (h[i] != i + (shared_mem ? 7 : 6))
+ abort ();
+ h[i] = i + 8;
+ }
+
+ {
+ int *d_ = (int *) acc_pcopyin (h, S);
+ if (d_ != d)
+ abort ();
+ }
+
+#pragma acc parallel loop deviceptr (d)
+ for (int i = 0; i < N; ++i)
+ {
+ if (d[i] != i + (shared_mem ? 8 : 7))
+ abort ();
+ d[i] = i + 9;
+ }
- d2 = acc_pcreate (h, N);
- if (!d2)
+ for (int i = 0; i < N; ++i)
+ {
+ if (h[i] != i + (shared_mem ? 9 : 8))
+ abort ();
+ h[i] = i + 10;
+ }
+
+ acc_copyout (h, S);
+ d = NULL;
+ if (!shared_mem)
+ if (acc_is_present (h, S))
+ abort ();
+
+ for (int i = 0; i < N; ++i)
+ {
+ if (h[i] != i + (shared_mem ? 10 : 9))
+ abort ();
+ }
+
+ d = (int *) acc_pcopyin (h, S);
+ if (!d)
+ abort ();
+ if (shared_mem)
+ if (h != d)
+ abort ();
+ if (!acc_is_present (h, S))
abort ();
- if (d1 != d2)
+#pragma acc parallel loop deviceptr (d)
+ for (int i = 0; i < N; ++i)
+ {
+ if (d[i] != i + (shared_mem ? 10 : 9))
+ abort ();
+ d[i] = i + 11;
+ }
+
+ for (int i = 0; i < N; ++i)
+ {
+ if (h[i] != i + (shared_mem ? 11 : 9))
+ abort ();
+ h[i] = i + 12;
+ }
+
+ {
+ int *d_ = (int *) acc_pcopyin (h, S);
+ if (d_ != d)
+ abort ();
+ }
+
+#pragma acc parallel loop deviceptr (d)
+ for (int i = 0; i < N; ++i)
+ {
+ if (d[i] != i + (shared_mem ? 12 : 11))
+ abort ();
+ d[i] = i + 13;
+ }
+
+ for (int i = 0; i < N; ++i)
+ {
+ if (h[i] != i + (shared_mem ? 13 : 12))
+ abort ();
+ h[i] = i + 14;
+ }
+
+ {
+ int *d_ = (int *) acc_pcreate (h, S);
+ if (d_ != d)
+ abort ();
+ }
+
+#pragma acc parallel loop deviceptr (d)
+ for (int i = 0; i < N; ++i)
+ {
+ if (d[i] != i + (shared_mem ? 14 : 13))
+ abort ();
+ d[i] = i + 15;
+ }
+
+ for (int i = 0; i < N; ++i)
+ {
+ if (h[i] != i + (shared_mem ? 15 : 14))
+ abort ();
+ h[i] = i + 16;
+ }
+
+ {
+ int *d_ = (int *) acc_pcreate (h, S);
+ if (d_ != d)
+ abort ();
+ }
+
+#pragma acc parallel loop deviceptr (d)
+ for (int i = 0; i < N; ++i)
+ {
+ if (d[i] != i + (shared_mem ? 16 : 15))
+ abort ();
+ d[i] = i + 17;
+ }
+
+ for (int i = 0; i < N; ++i)
+ {
+ if (h[i] != i + (shared_mem ? 17 : 16))
+ abort ();
+ h[i] = i + 18;
+ }
+
+ acc_update_self (h, S);
+ if (!acc_is_present (h, S))
abort ();
- acc_delete (h, N);
+ for (int i = 0; i < N; ++i)
+ {
+ if (h[i] != i + (shared_mem ? 18 : 17))
+ abort ();
+ }
+
+ acc_delete (h, S);
+ d = NULL;
+ if (!shared_mem)
+ if (acc_is_present (h, S))
+ abort();
free (h);
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/lib-38.c libgomp/testsuite/libgomp.oacc-c-c++-common/lib-38.c
deleted file mode 100644
index 05d8498..0000000
--- libgomp/testsuite/libgomp.oacc-c-c++-common/lib-38.c
+++ /dev/null
@@ -1,64 +0,0 @@
-/* { dg-do run } */
-/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
-
-#include <string.h>
-#include <stdlib.h>
-#include <openacc.h>
-
-int
-main (int argc, char **argv)
-{
- const int N = 256;
- int i;
- unsigned char *h;
- void *d1, *d2;
-
- h = (unsigned char *) malloc (N);
-
- for (i = 0; i < N; i++)
- {
- h[i] = i;
- }
-
- d1 = acc_present_or_copyin (h, N);
- if (!d1)
- abort ();
-
- for (i = 0; i < N; i++)
- {
- h[i] = 0xab;
- }
-
- d2 = acc_present_or_copyin (h, N);
- if (!d2)
- abort ();
-
- if (d1 != d2)
- abort ();
-
- memset (&h[0], 0, N);
-
- acc_copyout (h, N);
-
- for (i = 0; i < N; i++)
- {
- if (h[i] != i)
- abort ();
- }
-
- d2 = acc_pcopyin (h, N);
- if (!d2)
- abort ();
-
- acc_copyout (h, N);
-
- for (i = 0; i < N; i++)
- {
- if (h[i] != i)
- abort ();
- }
-
- free (h);
-
- return 0;
-}
Committed to gomp-4_0-branch in r248414:
commit 705bc833826559e64a7ccd821fcf419bfaf2b727
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Wed May 24 13:25:42 2017 +0000
C/C++ OpenACC: acc_pcopyin, acc_pcreate
libgomp/
* openacc.h (acc_pcopyin, acc_pcreate): Provide prototypes instead
of preprocessor definitions.
* libgomp.h (strong_alias): Guard by "#ifdef
HAVE_ATTRIBUTE_ALIAS".
* oacc-mem.c: Provide "acc_pcreate" as alias for
"acc_present_or_create", and "acc_pcopyin" as alias for
"acc_present_or_copyin".
* libgomp.map: New version "OACC_2.0.1".
(OACC_2.0.1): Add "acc_pcopyin", and "acc_pcreate".
* testsuite/libgomp.oacc-c-c++-common/lib-38.c: Remove, merging
its content into...
* testsuite/libgomp.oacc-c-c++-common/lib-32.c: ... this file.
Extend testing.
trunk r248410
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@248414 138bc75d-0d04-0410-961f-82ee72b054a4
---
libgomp/ChangeLog.gomp | 14 ++
libgomp/libgomp.h | 5 +-
libgomp/libgomp.map | 8 +-
libgomp/oacc-mem.c | 24 +-
libgomp/openacc.h | 7 +-
.../testsuite/libgomp.oacc-c-c++-common/lib-32.c | 241 +++++++++++++++++++--
.../testsuite/libgomp.oacc-c-c++-common/lib-38.c | 64 ------
7 files changed, 271 insertions(+), 92 deletions(-)
diff --git libgomp/ChangeLog.gomp libgomp/ChangeLog.gomp
index 15e0c58..4ab1004 100644
--- libgomp/ChangeLog.gomp
+++ libgomp/ChangeLog.gomp
@@ -1,5 +1,19 @@
2017-05-24 Thomas Schwinge <thomas@codesourcery.com>
+ * openacc.h (acc_pcopyin, acc_pcreate): Provide prototypes instead
+ of preprocessor definitions.
+ * libgomp.h (strong_alias): Guard by "#ifdef
+ HAVE_ATTRIBUTE_ALIAS".
+ * oacc-mem.c: Provide "acc_pcreate" as alias for
+ "acc_present_or_create", and "acc_pcopyin" as alias for
+ "acc_present_or_copyin".
+ * libgomp.map: New version "OACC_2.0.1".
+ (OACC_2.0.1): Add "acc_pcopyin", and "acc_pcreate".
+ * testsuite/libgomp.oacc-c-c++-common/lib-38.c: Remove, merging
+ its content into...
+ * testsuite/libgomp.oacc-c-c++-common/lib-32.c: ... this file.
+ Extend testing.
+
* plugin/plugin-nvptx.c (nvptx_get_num_devices): Debugging output
when disabling nvptx offloading.
diff --git libgomp/libgomp.h libgomp/libgomp.h
index a07a604..50916f2 100644
--- libgomp/libgomp.h
+++ libgomp/libgomp.h
@@ -1049,8 +1049,6 @@ extern void gomp_set_nest_lock_25 (omp_nest_lock_25_t *) __GOMP_NOTHROW;
extern void gomp_unset_nest_lock_25 (omp_nest_lock_25_t *) __GOMP_NOTHROW;
extern int gomp_test_nest_lock_25 (omp_nest_lock_25_t *) __GOMP_NOTHROW;
-# define strong_alias(fn, al) \
- extern __typeof (fn) al __attribute__ ((alias (#fn)));
# define omp_lock_symver(fn) \
__asm (".symver g" #fn "_30, " #fn "@@OMP_3.0"); \
__asm (".symver g" #fn "_25, " #fn "@OMP_1.0");
@@ -1074,6 +1072,9 @@ extern int gomp_test_nest_lock_25 (omp_nest_lock_25_t *) __GOMP_NOTHROW;
#endif
#ifdef HAVE_ATTRIBUTE_ALIAS
+# define strong_alias(fn, al) \
+ extern __typeof (fn) al __attribute__ ((alias (#fn)));
+
# define ialias_ulp ialias_str1(__USER_LABEL_PREFIX__)
# define ialias_str1(x) ialias_str2(x)
# define ialias_str2(x) #x
diff --git libgomp/libgomp.map libgomp/libgomp.map
index cee0bc9..2461c00 100644
--- libgomp/libgomp.map
+++ libgomp/libgomp.map
@@ -378,6 +378,12 @@ OACC_2.0 {
acc_set_cuda_stream;
};
+OACC_2.0.1 {
+ global:
+ acc_pcopyin;
+ acc_pcreate;
+} OACC_2.0;
+
OACC_2.5 {
global:
acc_copyin_async;
@@ -430,7 +436,7 @@ OACC_2.5 {
acc_update_self_async_32_h_;
acc_update_self_async_64_h_;
acc_update_self_async_array_h_;
-} OACC_2.0;
+} OACC_2.0.1;
GOACC_2.0 {
global:
diff --git libgomp/oacc-mem.c libgomp/oacc-mem.c
index 6a1b25d..59a6f93 100644
--- libgomp/oacc-mem.c
+++ libgomp/oacc-mem.c
@@ -654,15 +654,23 @@ 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
-extern void *acc_present_or_create (void *, size_t)
- __attribute__((alias ("acc_create")));
+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 *
@@ -678,15 +686,23 @@ 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
-extern void *acc_present_or_copyin (void *, size_t)
- __attribute__((alias ("acc_copyin")));
+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_copyin (h, s);
+}
#endif
#define FLAG_COPYOUT (1 << 0)
diff --git libgomp/openacc.h libgomp/openacc.h
index eb7c888..67d8639 100644
--- libgomp/openacc.h
+++ libgomp/openacc.h
@@ -94,8 +94,10 @@ void acc_free (void *) __GOACC_NOTHROW;
the standard specifies otherwise. */
void *acc_copyin (void *, size_t) __GOACC_NOTHROW;
void *acc_present_or_copyin (void *, size_t) __GOACC_NOTHROW;
+void *acc_pcopyin (void *, size_t) __GOACC_NOTHROW;
void *acc_create (void *, size_t) __GOACC_NOTHROW;
void *acc_present_or_create (void *, size_t) __GOACC_NOTHROW;
+void *acc_pcreate (void *, size_t) __GOACC_NOTHROW;
void acc_copyout (void *, size_t) __GOACC_NOTHROW;
void acc_delete (void *, size_t) __GOACC_NOTHROW;
void acc_update_device (void *, size_t) __GOACC_NOTHROW;
@@ -124,11 +126,6 @@ void acc_copyout_finalize_async (void *, size_t, int) __GOACC_NOTHROW;
void acc_delete_finalize (void *, size_t) __GOACC_NOTHROW;
void acc_delete_finalize_async (void *, size_t, int) __GOACC_NOTHROW;
-/* Old names. OpenACC does not specify whether these can or must
- not be macros, inlines or aliases for the new names. */
-#define acc_pcreate acc_present_or_create
-#define acc_pcopyin acc_present_or_copyin
-
/* CUDA-specific routines. */
void *acc_get_current_cuda_device (void) __GOACC_NOTHROW;
void *acc_get_current_cuda_context (void) __GOACC_NOTHROW;
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/lib-32.c libgomp/testsuite/libgomp.oacc-c-c++-common/lib-32.c
index e3f87a8..666d4b8 100644
--- libgomp/testsuite/libgomp.oacc-c-c++-common/lib-32.c
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/lib-32.c
@@ -1,36 +1,245 @@
-/* { dg-do run } */
+/* acc_present_or_create, acc_present_or_copyin, etc. */
+#include <stdbool.h>
#include <stdlib.h>
#include <openacc.h>
int
main (int argc, char **argv)
{
- const int N = 256;
- unsigned char *h;
- void *d1, *d2;
+ int *h, *d;
+ const int N = 10000;
+ const int S = N * sizeof *h;
+ bool shared_mem;
- h = (unsigned char *) malloc (N);
-
- d1 = acc_present_or_create (h, N);
- if (!d1)
+ h = (int *) malloc (S);
+ if (!h)
abort ();
+ for (int i = 0; i < N; ++i)
+ h[i] = i + 0;
- d2 = acc_present_or_create (h, N);
- if (!d2)
- abort ();
+ shared_mem = acc_is_present (h, S);
- if (d1 != d2)
+ d = (int *) acc_present_or_create (h, S);
+ if (!d)
abort ();
+ if (shared_mem)
+ if (h != d)
+ abort ();
+ if (!acc_is_present (h, S))
+ abort ();
+
+#pragma acc parallel loop deviceptr (d)
+ for (int i = 0; i < N; ++i)
+ {
+ d[i] = i + 1;
+ }
+
+ for (int i = 0; i < N; ++i)
+ {
+ if (h[i] != i + (shared_mem ? 1 : 0))
+ abort ();
+ h[i] = i + 2;
+ }
+
+ {
+ int *d_ = (int *) acc_present_or_create (h, S);
+ if (d_ != d)
+ abort ();
+ }
+
+#pragma acc parallel loop deviceptr (d)
+ for (int i = 0; i < N; ++i)
+ {
+ if (d[i] != i + (shared_mem ? 2 : 1))
+ abort ();
+ d[i] = i + 3;
+ }
+
+ for (int i = 0; i < N; ++i)
+ {
+ if (h[i] != i + (shared_mem ? 3 : 2))
+ abort ();
+ h[i] = i + 4;
+ }
+
+ {
+ int *d_ = (int *) acc_pcreate (h, S);
+ if (d_ != d)
+ abort ();
+ }
+
+#pragma acc parallel loop deviceptr (d)
+ for (int i = 0; i < N; ++i)
+ {
+ if (d[i] != i + (shared_mem ? 4 : 3))
+ abort ();
+ d[i] = i + 5;
+ }
+
+ for (int i = 0; i < N; ++i)
+ {
+ if (h[i] != i + (shared_mem ? 5 : 4))
+ abort ();
+ h[i] = i + 6;
+ }
+
+ {
+ int *d_ = (int *) acc_present_or_copyin (h, S);
+ if (d_ != d)
+ abort ();
+ }
+
+#pragma acc parallel loop deviceptr (d)
+ for (int i = 0; i < N; ++i)
+ {
+ if (d[i] != i + (shared_mem ? 6 : 5))
+ abort ();
+ d[i] = i + 7;
+ }
+
+ for (int i = 0; i < N; ++i)
+ {
+ if (h[i] != i + (shared_mem ? 7 : 6))
+ abort ();
+ h[i] = i + 8;
+ }
+
+ {
+ int *d_ = (int *) acc_pcopyin (h, S);
+ if (d_ != d)
+ abort ();
+ }
+
+#pragma acc parallel loop deviceptr (d)
+ for (int i = 0; i < N; ++i)
+ {
+ if (d[i] != i + (shared_mem ? 8 : 7))
+ abort ();
+ d[i] = i + 9;
+ }
- d2 = acc_pcreate (h, N);
- if (!d2)
+ for (int i = 0; i < N; ++i)
+ {
+ if (h[i] != i + (shared_mem ? 9 : 8))
+ abort ();
+ h[i] = i + 10;
+ }
+
+ acc_copyout_finalize (h, S);
+ d = NULL;
+ if (!shared_mem)
+ if (acc_is_present (h, S))
+ abort ();
+
+ for (int i = 0; i < N; ++i)
+ {
+ if (h[i] != i + (shared_mem ? 10 : 9))
+ abort ();
+ }
+
+ d = (int *) acc_pcopyin (h, S);
+ if (!d)
+ abort ();
+ if (shared_mem)
+ if (h != d)
+ abort ();
+ if (!acc_is_present (h, S))
abort ();
- if (d1 != d2)
+#pragma acc parallel loop deviceptr (d)
+ for (int i = 0; i < N; ++i)
+ {
+ if (d[i] != i + (shared_mem ? 10 : 9))
+ abort ();
+ d[i] = i + 11;
+ }
+
+ for (int i = 0; i < N; ++i)
+ {
+ if (h[i] != i + (shared_mem ? 11 : 9))
+ abort ();
+ h[i] = i + 12;
+ }
+
+ {
+ int *d_ = (int *) acc_pcopyin (h, S);
+ if (d_ != d)
+ abort ();
+ }
+
+#pragma acc parallel loop deviceptr (d)
+ for (int i = 0; i < N; ++i)
+ {
+ if (d[i] != i + (shared_mem ? 12 : 11))
+ abort ();
+ d[i] = i + 13;
+ }
+
+ for (int i = 0; i < N; ++i)
+ {
+ if (h[i] != i + (shared_mem ? 13 : 12))
+ abort ();
+ h[i] = i + 14;
+ }
+
+ {
+ int *d_ = (int *) acc_pcreate (h, S);
+ if (d_ != d)
+ abort ();
+ }
+
+#pragma acc parallel loop deviceptr (d)
+ for (int i = 0; i < N; ++i)
+ {
+ if (d[i] != i + (shared_mem ? 14 : 13))
+ abort ();
+ d[i] = i + 15;
+ }
+
+ for (int i = 0; i < N; ++i)
+ {
+ if (h[i] != i + (shared_mem ? 15 : 14))
+ abort ();
+ h[i] = i + 16;
+ }
+
+ {
+ int *d_ = (int *) acc_pcreate (h, S);
+ if (d_ != d)
+ abort ();
+ }
+
+#pragma acc parallel loop deviceptr (d)
+ for (int i = 0; i < N; ++i)
+ {
+ if (d[i] != i + (shared_mem ? 16 : 15))
+ abort ();
+ d[i] = i + 17;
+ }
+
+ for (int i = 0; i < N; ++i)
+ {
+ if (h[i] != i + (shared_mem ? 17 : 16))
+ abort ();
+ h[i] = i + 18;
+ }
+
+ acc_update_self (h, S);
+ if (!acc_is_present (h, S))
abort ();
- acc_delete (h, N);
+ for (int i = 0; i < N; ++i)
+ {
+ if (h[i] != i + (shared_mem ? 18 : 17))
+ abort ();
+ }
+
+ acc_delete_finalize (h, S);
+ d = NULL;
+ if (!shared_mem)
+ if (acc_is_present (h, S))
+ abort();
free (h);
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/lib-38.c libgomp/testsuite/libgomp.oacc-c-c++-common/lib-38.c
deleted file mode 100644
index 1756867..0000000
--- libgomp/testsuite/libgomp.oacc-c-c++-common/lib-38.c
+++ /dev/null
@@ -1,64 +0,0 @@
-/* { dg-do run } */
-/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
-
-#include <string.h>
-#include <stdlib.h>
-#include <openacc.h>
-
-int
-main (int argc, char **argv)
-{
- const int N = 256;
- int i;
- unsigned char *h;
- void *d1, *d2;
-
- h = (unsigned char *) malloc (N);
-
- for (i = 0; i < N; i++)
- {
- h[i] = i;
- }
-
- d1 = acc_present_or_copyin (h, N);
- if (!d1)
- abort ();
-
- for (i = 0; i < N; i++)
- {
- h[i] = 0xab;
- }
-
- d2 = acc_present_or_copyin (h, N);
- if (!d2)
- abort ();
-
- if (d1 != d2)
- abort ();
-
- memset (&h[0], 0, N);
-
- acc_copyout_finalize (h, N);
-
- for (i = 0; i < N; i++)
- {
- if (h[i] != i)
- abort ();
- }
-
- d2 = acc_pcopyin (h, N);
- if (!d2)
- abort ();
-
- acc_copyout (h, N);
-
- for (i = 0; i < N; i++)
- {
- if (h[i] != i)
- abort ();
- }
-
- free (h);
-
- return 0;
-}
Grüße
Thomas
^ permalink raw reply [flat|nested] 8+ messages in thread
* Re: Translate libgomp.oacc-c-c++-common/lib-32.c into Fortran
2017-05-23 11:12 ` Jakub Jelinek
@ 2017-05-24 13:45 ` Thomas Schwinge
0 siblings, 0 replies; 8+ messages in thread
From: Thomas Schwinge @ 2017-05-24 13:45 UTC (permalink / raw)
To: Jakub Jelinek, gcc-patches, fortran
Hi!
On Tue, 23 May 2017 13:08:22 +0200, Jakub Jelinek <jakub@redhat.com> wrote:
> On Mon, May 22, 2017 at 04:29:51PM +0200, Thomas Schwinge wrote:
> > On Mon, 22 May 2017 16:26:48 +0200, I wrote:
> > > C/C++ OpenACC: acc_pcopyin, acc_pcreate
> >
> > > libgomp/
> > > [...]
> > > * testsuite/libgomp.oacc-c-c++-common/lib-38.c: Remove, merging
> > > its content into...
> > > * testsuite/libgomp.oacc-c-c++-common/lib-32.c: ... this file.
> > > Extend testing.
> >
> > ... which I then translated into two Fortran variants; OK for trunk?
> Ok.
Thanks. As posted, committed to trunk in r248411:
commit 4a5596a04ac52ae4df1347e7980e0d60fcd1e84b
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Wed May 24 13:23:45 2017 +0000
Translate libgomp.oacc-c-c++-common/lib-32.c into Fortran
libgomp/
* testsuite/libgomp.oacc-fortran/lib-32-1.f: New file.
* testsuite/libgomp.oacc-fortran/lib-32-2.f: Likewise.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@248411 138bc75d-0d04-0410-961f-82ee72b054a4
---
libgomp/ChangeLog | 3 +
.../testsuite/libgomp.oacc-c-c++-common/lib-32.c | 1 +
libgomp/testsuite/libgomp.oacc-fortran/lib-32-1.f | 179 +++++++++++++++++++++
libgomp/testsuite/libgomp.oacc-fortran/lib-32-2.f | 173 ++++++++++++++++++++
4 files changed, 356 insertions(+)
diff --git libgomp/ChangeLog libgomp/ChangeLog
index 4a95755..af8c5ca 100644
--- libgomp/ChangeLog
+++ libgomp/ChangeLog
@@ -1,5 +1,8 @@
2017-05-24 Thomas Schwinge <thomas@codesourcery.com>
+ * testsuite/libgomp.oacc-fortran/lib-32-1.f: New file.
+ * testsuite/libgomp.oacc-fortran/lib-32-2.f: Likewise.
+
* openacc.h (acc_pcopyin, acc_pcreate): Provide prototypes instead
of preprocessor definitions.
* libgomp.h (strong_alias): Guard by "#ifdef
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/lib-32.c libgomp/testsuite/libgomp.oacc-c-c++-common/lib-32.c
index 6a9e995..1696fb6 100644
--- libgomp/testsuite/libgomp.oacc-c-c++-common/lib-32.c
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/lib-32.c
@@ -1,4 +1,5 @@
/* acc_present_or_create, acc_present_or_copyin, etc. */
+/* See also Fortran variants in "../libgomp.oacc-fortran/lib-32*". */
#include <stdbool.h>
#include <stdlib.h>
diff --git libgomp/testsuite/libgomp.oacc-fortran/lib-32-1.f libgomp/testsuite/libgomp.oacc-fortran/lib-32-1.f
new file mode 100644
index 0000000..4606d77
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-fortran/lib-32-1.f
@@ -0,0 +1,179 @@
+! ACC_PRESENT_OR_CREATE, ACC_PRESENT_OR_COPYIN, etc.
+! Variant of "../libgomp.oacc-c-c++-common/lib-32.c".
+! Variant using "openacc_lib.h".
+
+! { dg-do run }
+
+ PROGRAM MAIN
+ IMPLICIT NONE
+ INCLUDE "openacc_lib.h"
+
+ INTEGER, PARAMETER :: N = 10000
+ INTEGER, ALLOCATABLE :: H(:)
+ INTEGER :: I
+ LOGICAL :: SHARED_MEM
+
+ ALLOCATE (H(N))
+ DO I = 1, N
+ H(I) = I + 0
+ END DO
+
+ SHARED_MEM = ACC_IS_PRESENT (H)
+
+ CALL ACC_PRESENT_OR_CREATE (H)
+ IF (.NOT. ACC_IS_PRESENT (H)) CALL ABORT
+
+!$ACC PARALLEL LOOP DEFAULT (PRESENT)
+ DO I = 1, N
+ H(I) = I + 1
+ END DO
+!$ACC END PARALLEL LOOP
+
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (1, 0, SHARED_MEM)) CALL ABORT
+ H(I) = I + 2
+ END DO
+
+ CALL ACC_PRESENT_OR_CREATE (H)
+
+!$ACC PARALLEL LOOP DEFAULT (PRESENT)
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (2, 1, SHARED_MEM)) CALL ABORT
+ H(I) = I + 3
+ END DO
+!$ACC END PARALLEL LOOP
+
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (3, 2, SHARED_MEM)) CALL ABORT
+ H(I) = I + 4
+ END DO
+
+! CALL ACC_PCREATE (H)
+ CALL ACC_PRESENT_OR_CREATE (H)
+
+!$ACC PARALLEL LOOP DEFAULT (PRESENT)
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (4, 3, SHARED_MEM)) CALL ABORT
+ H(I) = I + 5
+ END DO
+!$ACC END PARALLEL LOOP
+
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (5, 4, SHARED_MEM)) CALL ABORT
+ H(I) = I + 6
+ END DO
+
+ CALL ACC_PRESENT_OR_COPYIN (H)
+
+!$ACC PARALLEL LOOP DEFAULT (PRESENT)
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (6, 5, SHARED_MEM)) CALL ABORT
+ H(I) = I + 7
+ END DO
+!$ACC END PARALLEL LOOP
+
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (7, 6, SHARED_MEM)) CALL ABORT
+ H(I) = I + 8
+ END DO
+
+! CALL ACC_PCOPYIN (H)
+ CALL ACC_PRESENT_OR_COPYIN (H)
+
+!$ACC PARALLEL LOOP DEFAULT (PRESENT)
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (8, 7, SHARED_MEM)) CALL ABORT
+ H(I) = I + 9
+ END DO
+!$ACC END PARALLEL LOOP
+
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (9, 8, SHARED_MEM)) CALL ABORT
+ H(I) = I + 10
+ END DO
+
+ CALL ACC_COPYOUT (H)
+ IF (.NOT. SHARED_MEM) THEN
+ IF (ACC_IS_PRESENT (H)) CALL ABORT
+ ENDIF
+
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (10, 9, SHARED_MEM)) CALL ABORT
+ END DO
+
+! CALL ACC_PCOPYIN (H)
+ CALL ACC_PRESENT_OR_COPYIN (H)
+ IF (.NOT. ACC_IS_PRESENT (H)) CALL ABORT
+
+!$ACC PARALLEL LOOP DEFAULT (PRESENT)
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (10, 9, SHARED_MEM)) CALL ABORT
+ H(I) = I + 11
+ END DO
+!$ACC END PARALLEL LOOP
+
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (11, 9, SHARED_MEM)) CALL ABORT
+ H(I) = I + 12
+ END DO
+
+! CALL ACC_PCOPYIN (H)
+ CALL ACC_PRESENT_OR_COPYIN (H)
+
+!$ACC PARALLEL LOOP DEFAULT (PRESENT)
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (12, 11, SHARED_MEM)) CALL ABORT
+ H(I) = I + 13
+ END DO
+!$ACC END PARALLEL LOOP
+
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (13, 12, SHARED_MEM)) CALL ABORT
+ H(I) = I + 14
+ END DO
+
+! CALL ACC_PCREATE (H)
+ CALL ACC_PRESENT_OR_CREATE (H)
+
+!$ACC PARALLEL LOOP DEFAULT (PRESENT)
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (14, 13, SHARED_MEM)) CALL ABORT
+ H(I) = I + 15
+ END DO
+!$ACC END PARALLEL LOOP
+
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (15, 14, SHARED_MEM)) CALL ABORT
+ H(I) = I + 16
+ END DO
+
+! CALL ACC_PCREATE (H)
+ CALL ACC_PRESENT_OR_CREATE (H)
+
+!$ACC PARALLEL LOOP DEFAULT (PRESENT)
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (16, 15, SHARED_MEM)) CALL ABORT
+ H(I) = I + 17
+ END DO
+!$ACC END PARALLEL LOOP
+
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (17, 16, SHARED_MEM)) CALL ABORT
+ H(I) = I + 18
+ END DO
+
+ CALL ACC_UPDATE_SELF (H)
+ IF (.NOT. ACC_IS_PRESENT (H)) CALL ABORT
+
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (18, 17, SHARED_MEM)) CALL ABORT
+ END DO
+
+ CALL ACC_DELETE (H)
+ IF (.NOT. SHARED_MEM) THEN
+ IF (ACC_IS_PRESENT (H)) CALL ABORT
+ ENDIF
+
+ DEALLOCATE (H)
+
+ END PROGRAM MAIN
diff --git libgomp/testsuite/libgomp.oacc-fortran/lib-32-2.f libgomp/testsuite/libgomp.oacc-fortran/lib-32-2.f
new file mode 100644
index 0000000..88f5566
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-fortran/lib-32-2.f
@@ -0,0 +1,173 @@
+! ACC_PRESENT_OR_CREATE, ACC_PRESENT_OR_COPYIN, etc.
+! Variant of "../libgomp.oacc-c-c++-common/lib-32.c".
+! Variant using the "openacc" module.
+
+! { dg-do run }
+
+ PROGRAM MAIN
+ USE OPENACC
+ IMPLICIT NONE
+
+ INTEGER, PARAMETER :: N = 10000
+ INTEGER, ALLOCATABLE :: H(:)
+ INTEGER :: I
+ LOGICAL :: SHARED_MEM
+
+ ALLOCATE (H(N))
+ DO I = 1, N
+ H(I) = I + 0
+ END DO
+
+ SHARED_MEM = ACC_IS_PRESENT (H)
+
+ CALL ACC_PRESENT_OR_CREATE (H)
+ IF (.NOT. ACC_IS_PRESENT (H)) CALL ABORT
+
+!$ACC PARALLEL LOOP DEFAULT (PRESENT)
+ DO I = 1, N
+ H(I) = I + 1
+ END DO
+!$ACC END PARALLEL LOOP
+
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (1, 0, SHARED_MEM)) CALL ABORT
+ H(I) = I + 2
+ END DO
+
+ CALL ACC_PRESENT_OR_CREATE (H)
+
+!$ACC PARALLEL LOOP DEFAULT (PRESENT)
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (2, 1, SHARED_MEM)) CALL ABORT
+ H(I) = I + 3
+ END DO
+!$ACC END PARALLEL LOOP
+
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (3, 2, SHARED_MEM)) CALL ABORT
+ H(I) = I + 4
+ END DO
+
+ CALL ACC_PCREATE (H)
+
+!$ACC PARALLEL LOOP DEFAULT (PRESENT)
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (4, 3, SHARED_MEM)) CALL ABORT
+ H(I) = I + 5
+ END DO
+!$ACC END PARALLEL LOOP
+
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (5, 4, SHARED_MEM)) CALL ABORT
+ H(I) = I + 6
+ END DO
+
+ CALL ACC_PRESENT_OR_COPYIN (H)
+
+!$ACC PARALLEL LOOP DEFAULT (PRESENT)
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (6, 5, SHARED_MEM)) CALL ABORT
+ H(I) = I + 7
+ END DO
+!$ACC END PARALLEL LOOP
+
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (7, 6, SHARED_MEM)) CALL ABORT
+ H(I) = I + 8
+ END DO
+
+ CALL ACC_PCOPYIN (H)
+
+!$ACC PARALLEL LOOP DEFAULT (PRESENT)
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (8, 7, SHARED_MEM)) CALL ABORT
+ H(I) = I + 9
+ END DO
+!$ACC END PARALLEL LOOP
+
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (9, 8, SHARED_MEM)) CALL ABORT
+ H(I) = I + 10
+ END DO
+
+ CALL ACC_COPYOUT (H)
+ IF (.NOT. SHARED_MEM) THEN
+ IF (ACC_IS_PRESENT (H)) CALL ABORT
+ ENDIF
+
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (10, 9, SHARED_MEM)) CALL ABORT
+ END DO
+
+ CALL ACC_PCOPYIN (H)
+ IF (.NOT. ACC_IS_PRESENT (H)) CALL ABORT
+
+!$ACC PARALLEL LOOP DEFAULT (PRESENT)
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (10, 9, SHARED_MEM)) CALL ABORT
+ H(I) = I + 11
+ END DO
+!$ACC END PARALLEL LOOP
+
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (11, 9, SHARED_MEM)) CALL ABORT
+ H(I) = I + 12
+ END DO
+
+ CALL ACC_PCOPYIN (H)
+
+!$ACC PARALLEL LOOP DEFAULT (PRESENT)
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (12, 11, SHARED_MEM)) CALL ABORT
+ H(I) = I + 13
+ END DO
+!$ACC END PARALLEL LOOP
+
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (13, 12, SHARED_MEM)) CALL ABORT
+ H(I) = I + 14
+ END DO
+
+ CALL ACC_PCREATE (H)
+
+!$ACC PARALLEL LOOP DEFAULT (PRESENT)
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (14, 13, SHARED_MEM)) CALL ABORT
+ H(I) = I + 15
+ END DO
+!$ACC END PARALLEL LOOP
+
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (15, 14, SHARED_MEM)) CALL ABORT
+ H(I) = I + 16
+ END DO
+
+ CALL ACC_PCREATE (H)
+
+!$ACC PARALLEL LOOP DEFAULT (PRESENT)
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (16, 15, SHARED_MEM)) CALL ABORT
+ H(I) = I + 17
+ END DO
+!$ACC END PARALLEL LOOP
+
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (17, 16, SHARED_MEM)) CALL ABORT
+ H(I) = I + 18
+ END DO
+
+ CALL ACC_UPDATE_SELF (H)
+ IF (.NOT. ACC_IS_PRESENT (H)) CALL ABORT
+
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (18, 17, SHARED_MEM)) CALL ABORT
+ END DO
+
+ CALL ACC_DELETE (H)
+ IF (.NOT. SHARED_MEM) THEN
+ IF (ACC_IS_PRESENT (H)) CALL ABORT
+ ENDIF
+
+ DEALLOCATE (H)
+
+ END PROGRAM MAIN
Committed to gomp-4_0-branch in r248415:
commit 7621601fad34df2509e572e961ea7dae20d7b225
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Wed May 24 13:25:54 2017 +0000
Translate libgomp.oacc-c-c++-common/lib-32.c into Fortran
libgomp/
* testsuite/libgomp.oacc-fortran/lib-32-1.f: New file.
* testsuite/libgomp.oacc-fortran/lib-32-2.f: Likewise.
trunk r248411
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@248415 138bc75d-0d04-0410-961f-82ee72b054a4
---
libgomp/ChangeLog.gomp | 3 +
.../testsuite/libgomp.oacc-c-c++-common/lib-32.c | 1 +
libgomp/testsuite/libgomp.oacc-fortran/lib-32-1.f | 179 +++++++++++++++++++++
libgomp/testsuite/libgomp.oacc-fortran/lib-32-2.f | 173 ++++++++++++++++++++
4 files changed, 356 insertions(+)
diff --git libgomp/ChangeLog.gomp libgomp/ChangeLog.gomp
index 4ab1004..de4eeb4 100644
--- libgomp/ChangeLog.gomp
+++ libgomp/ChangeLog.gomp
@@ -1,5 +1,8 @@
2017-05-24 Thomas Schwinge <thomas@codesourcery.com>
+ * testsuite/libgomp.oacc-fortran/lib-32-1.f: New file.
+ * testsuite/libgomp.oacc-fortran/lib-32-2.f: Likewise.
+
* openacc.h (acc_pcopyin, acc_pcreate): Provide prototypes instead
of preprocessor definitions.
* libgomp.h (strong_alias): Guard by "#ifdef
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/lib-32.c libgomp/testsuite/libgomp.oacc-c-c++-common/lib-32.c
index 666d4b8..9ec3453 100644
--- libgomp/testsuite/libgomp.oacc-c-c++-common/lib-32.c
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/lib-32.c
@@ -1,4 +1,5 @@
/* acc_present_or_create, acc_present_or_copyin, etc. */
+/* See also Fortran variants in "../libgomp.oacc-fortran/lib-32*". */
#include <stdbool.h>
#include <stdlib.h>
diff --git libgomp/testsuite/libgomp.oacc-fortran/lib-32-1.f libgomp/testsuite/libgomp.oacc-fortran/lib-32-1.f
new file mode 100644
index 0000000..4ed866b
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-fortran/lib-32-1.f
@@ -0,0 +1,179 @@
+! ACC_PRESENT_OR_CREATE, ACC_PRESENT_OR_COPYIN, etc.
+! Variant of "../libgomp.oacc-c-c++-common/lib-32.c".
+! Variant using "openacc_lib.h".
+
+! { dg-do run }
+
+ PROGRAM MAIN
+ IMPLICIT NONE
+ INCLUDE "openacc_lib.h"
+
+ INTEGER, PARAMETER :: N = 10000
+ INTEGER, ALLOCATABLE :: H(:)
+ INTEGER :: I
+ LOGICAL :: SHARED_MEM
+
+ ALLOCATE (H(N))
+ DO I = 1, N
+ H(I) = I + 0
+ END DO
+
+ SHARED_MEM = ACC_IS_PRESENT (H)
+
+ CALL ACC_PRESENT_OR_CREATE (H)
+ IF (.NOT. ACC_IS_PRESENT (H)) CALL ABORT
+
+!$ACC PARALLEL LOOP DEFAULT (PRESENT)
+ DO I = 1, N
+ H(I) = I + 1
+ END DO
+!$ACC END PARALLEL LOOP
+
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (1, 0, SHARED_MEM)) CALL ABORT
+ H(I) = I + 2
+ END DO
+
+ CALL ACC_PRESENT_OR_CREATE (H)
+
+!$ACC PARALLEL LOOP DEFAULT (PRESENT)
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (2, 1, SHARED_MEM)) CALL ABORT
+ H(I) = I + 3
+ END DO
+!$ACC END PARALLEL LOOP
+
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (3, 2, SHARED_MEM)) CALL ABORT
+ H(I) = I + 4
+ END DO
+
+! CALL ACC_PCREATE (H)
+ CALL ACC_PRESENT_OR_CREATE (H)
+
+!$ACC PARALLEL LOOP DEFAULT (PRESENT)
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (4, 3, SHARED_MEM)) CALL ABORT
+ H(I) = I + 5
+ END DO
+!$ACC END PARALLEL LOOP
+
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (5, 4, SHARED_MEM)) CALL ABORT
+ H(I) = I + 6
+ END DO
+
+ CALL ACC_PRESENT_OR_COPYIN (H)
+
+!$ACC PARALLEL LOOP DEFAULT (PRESENT)
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (6, 5, SHARED_MEM)) CALL ABORT
+ H(I) = I + 7
+ END DO
+!$ACC END PARALLEL LOOP
+
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (7, 6, SHARED_MEM)) CALL ABORT
+ H(I) = I + 8
+ END DO
+
+! CALL ACC_PCOPYIN (H)
+ CALL ACC_PRESENT_OR_COPYIN (H)
+
+!$ACC PARALLEL LOOP DEFAULT (PRESENT)
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (8, 7, SHARED_MEM)) CALL ABORT
+ H(I) = I + 9
+ END DO
+!$ACC END PARALLEL LOOP
+
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (9, 8, SHARED_MEM)) CALL ABORT
+ H(I) = I + 10
+ END DO
+
+ CALL ACC_COPYOUT_FINALIZE (H)
+ IF (.NOT. SHARED_MEM) THEN
+ IF (ACC_IS_PRESENT (H)) CALL ABORT
+ ENDIF
+
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (10, 9, SHARED_MEM)) CALL ABORT
+ END DO
+
+! CALL ACC_PCOPYIN (H)
+ CALL ACC_PRESENT_OR_COPYIN (H)
+ IF (.NOT. ACC_IS_PRESENT (H)) CALL ABORT
+
+!$ACC PARALLEL LOOP DEFAULT (PRESENT)
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (10, 9, SHARED_MEM)) CALL ABORT
+ H(I) = I + 11
+ END DO
+!$ACC END PARALLEL LOOP
+
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (11, 9, SHARED_MEM)) CALL ABORT
+ H(I) = I + 12
+ END DO
+
+! CALL ACC_PCOPYIN (H)
+ CALL ACC_PRESENT_OR_COPYIN (H)
+
+!$ACC PARALLEL LOOP DEFAULT (PRESENT)
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (12, 11, SHARED_MEM)) CALL ABORT
+ H(I) = I + 13
+ END DO
+!$ACC END PARALLEL LOOP
+
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (13, 12, SHARED_MEM)) CALL ABORT
+ H(I) = I + 14
+ END DO
+
+! CALL ACC_PCREATE (H)
+ CALL ACC_PRESENT_OR_CREATE (H)
+
+!$ACC PARALLEL LOOP DEFAULT (PRESENT)
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (14, 13, SHARED_MEM)) CALL ABORT
+ H(I) = I + 15
+ END DO
+!$ACC END PARALLEL LOOP
+
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (15, 14, SHARED_MEM)) CALL ABORT
+ H(I) = I + 16
+ END DO
+
+! CALL ACC_PCREATE (H)
+ CALL ACC_PRESENT_OR_CREATE (H)
+
+!$ACC PARALLEL LOOP DEFAULT (PRESENT)
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (16, 15, SHARED_MEM)) CALL ABORT
+ H(I) = I + 17
+ END DO
+!$ACC END PARALLEL LOOP
+
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (17, 16, SHARED_MEM)) CALL ABORT
+ H(I) = I + 18
+ END DO
+
+ CALL ACC_UPDATE_SELF (H)
+ IF (.NOT. ACC_IS_PRESENT (H)) CALL ABORT
+
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (18, 17, SHARED_MEM)) CALL ABORT
+ END DO
+
+ CALL ACC_DELETE_FINALIZE (H)
+ IF (.NOT. SHARED_MEM) THEN
+ IF (ACC_IS_PRESENT (H)) CALL ABORT
+ ENDIF
+
+ DEALLOCATE (H)
+
+ END PROGRAM MAIN
diff --git libgomp/testsuite/libgomp.oacc-fortran/lib-32-2.f libgomp/testsuite/libgomp.oacc-fortran/lib-32-2.f
new file mode 100644
index 0000000..3e88ce2
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-fortran/lib-32-2.f
@@ -0,0 +1,173 @@
+! ACC_PRESENT_OR_CREATE, ACC_PRESENT_OR_COPYIN, etc.
+! Variant of "../libgomp.oacc-c-c++-common/lib-32.c".
+! Variant using the "openacc" module.
+
+! { dg-do run }
+
+ PROGRAM MAIN
+ USE OPENACC
+ IMPLICIT NONE
+
+ INTEGER, PARAMETER :: N = 10000
+ INTEGER, ALLOCATABLE :: H(:)
+ INTEGER :: I
+ LOGICAL :: SHARED_MEM
+
+ ALLOCATE (H(N))
+ DO I = 1, N
+ H(I) = I + 0
+ END DO
+
+ SHARED_MEM = ACC_IS_PRESENT (H)
+
+ CALL ACC_PRESENT_OR_CREATE (H)
+ IF (.NOT. ACC_IS_PRESENT (H)) CALL ABORT
+
+!$ACC PARALLEL LOOP DEFAULT (PRESENT)
+ DO I = 1, N
+ H(I) = I + 1
+ END DO
+!$ACC END PARALLEL LOOP
+
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (1, 0, SHARED_MEM)) CALL ABORT
+ H(I) = I + 2
+ END DO
+
+ CALL ACC_PRESENT_OR_CREATE (H)
+
+!$ACC PARALLEL LOOP DEFAULT (PRESENT)
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (2, 1, SHARED_MEM)) CALL ABORT
+ H(I) = I + 3
+ END DO
+!$ACC END PARALLEL LOOP
+
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (3, 2, SHARED_MEM)) CALL ABORT
+ H(I) = I + 4
+ END DO
+
+ CALL ACC_PCREATE (H)
+
+!$ACC PARALLEL LOOP DEFAULT (PRESENT)
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (4, 3, SHARED_MEM)) CALL ABORT
+ H(I) = I + 5
+ END DO
+!$ACC END PARALLEL LOOP
+
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (5, 4, SHARED_MEM)) CALL ABORT
+ H(I) = I + 6
+ END DO
+
+ CALL ACC_PRESENT_OR_COPYIN (H)
+
+!$ACC PARALLEL LOOP DEFAULT (PRESENT)
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (6, 5, SHARED_MEM)) CALL ABORT
+ H(I) = I + 7
+ END DO
+!$ACC END PARALLEL LOOP
+
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (7, 6, SHARED_MEM)) CALL ABORT
+ H(I) = I + 8
+ END DO
+
+ CALL ACC_PCOPYIN (H)
+
+!$ACC PARALLEL LOOP DEFAULT (PRESENT)
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (8, 7, SHARED_MEM)) CALL ABORT
+ H(I) = I + 9
+ END DO
+!$ACC END PARALLEL LOOP
+
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (9, 8, SHARED_MEM)) CALL ABORT
+ H(I) = I + 10
+ END DO
+
+ CALL ACC_COPYOUT_FINALIZE (H)
+ IF (.NOT. SHARED_MEM) THEN
+ IF (ACC_IS_PRESENT (H)) CALL ABORT
+ ENDIF
+
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (10, 9, SHARED_MEM)) CALL ABORT
+ END DO
+
+ CALL ACC_PCOPYIN (H)
+ IF (.NOT. ACC_IS_PRESENT (H)) CALL ABORT
+
+!$ACC PARALLEL LOOP DEFAULT (PRESENT)
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (10, 9, SHARED_MEM)) CALL ABORT
+ H(I) = I + 11
+ END DO
+!$ACC END PARALLEL LOOP
+
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (11, 9, SHARED_MEM)) CALL ABORT
+ H(I) = I + 12
+ END DO
+
+ CALL ACC_PCOPYIN (H)
+
+!$ACC PARALLEL LOOP DEFAULT (PRESENT)
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (12, 11, SHARED_MEM)) CALL ABORT
+ H(I) = I + 13
+ END DO
+!$ACC END PARALLEL LOOP
+
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (13, 12, SHARED_MEM)) CALL ABORT
+ H(I) = I + 14
+ END DO
+
+ CALL ACC_PCREATE (H)
+
+!$ACC PARALLEL LOOP DEFAULT (PRESENT)
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (14, 13, SHARED_MEM)) CALL ABORT
+ H(I) = I + 15
+ END DO
+!$ACC END PARALLEL LOOP
+
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (15, 14, SHARED_MEM)) CALL ABORT
+ H(I) = I + 16
+ END DO
+
+ CALL ACC_PCREATE (H)
+
+!$ACC PARALLEL LOOP DEFAULT (PRESENT)
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (16, 15, SHARED_MEM)) CALL ABORT
+ H(I) = I + 17
+ END DO
+!$ACC END PARALLEL LOOP
+
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (17, 16, SHARED_MEM)) CALL ABORT
+ H(I) = I + 18
+ END DO
+
+ CALL ACC_UPDATE_SELF (H)
+ IF (.NOT. ACC_IS_PRESENT (H)) CALL ABORT
+
+ DO I = 1, N
+ IF (H(I) .NE. I + MERGE (18, 17, SHARED_MEM)) CALL ABORT
+ END DO
+
+ CALL ACC_DELETE_FINALIZE (H)
+ IF (.NOT. SHARED_MEM) THEN
+ IF (ACC_IS_PRESENT (H)) CALL ABORT
+ ENDIF
+
+ DEALLOCATE (H)
+
+ END PROGRAM MAIN
Grüße
Thomas
^ permalink raw reply [flat|nested] 8+ messages in thread
end of thread, other threads:[~2017-05-24 13:43 UTC | newest]
Thread overview: 8+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2017-05-22 14:28 C/C++ OpenACC: acc_pcopyin, acc_pcreate Thomas Schwinge
2017-05-22 15:23 ` Translate libgomp.oacc-c-c++-common/lib-32.c into Fortran (was: C/C++ OpenACC: acc_pcopyin, acc_pcreate) Thomas Schwinge
2017-05-23 11:12 ` Jakub Jelinek
2017-05-24 13:45 ` Translate libgomp.oacc-c-c++-common/lib-32.c into Fortran Thomas Schwinge
2017-05-23 11:08 ` C/C++ OpenACC: acc_pcopyin, acc_pcreate Jakub Jelinek
2017-05-24 10:57 ` Thomas Schwinge
2017-05-24 11:24 ` Jakub Jelinek
2017-05-24 13:41 ` Thomas Schwinge
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).