Re: C/C++ OpenACC: acc_pcopyin, acc_pcreate

[Thomas Schwinge <thomas@codesourcery.com>]

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