public inbox for gcc-patches@gcc.gnu.org
 help / color / mirror / Atom feed
* 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
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).