Hi Jakub!
On Tue, 23 May 2017 13:07:15 +0200, Jakub Jelinek <ja...@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 <tho...@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
