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

7 files changed, 272 insertions, 87 deletions

diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index 14e95ef9ba9..4a9575567b3 100644
--- a/libgomp/ChangeLog
+++ b/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 a/libgomp/libgomp.h b/libgomp/libgomp.h
index 1769a489fe8..940b5b83f70 100644
--- a/libgomp/libgomp.h
+++ b/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 a/libgomp/libgomp.map b/libgomp/libgomp.map
index 4d42c42f441..b43c6deb851 100644
--- a/libgomp/libgomp.map
+++ b/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 a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c
index 2df220201ba..ff3ed49a586 100644
--- a/libgomp/oacc-mem.c
+++ b/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 a/libgomp/openacc.h b/libgomp/openacc.h
index 53d0c3955c1..ebccb1856d7 100644
--- a/libgomp/openacc.h
+++ b/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 a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-32.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-32.c
index e3f87a85672..6a9e995530f 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-32.c
+++ b/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 a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-38.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-38.c
deleted file mode 100644
index 05d8498c1f9..00000000000
--- a/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;
-}