C/C++ OpenACC: acc_pcopyin, acc_pcreate

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;
-}