public inbox for gcc-patches@gcc.gnu.org
 help / color / mirror / Atom feed
From: Thomas Schwinge <thomas@codesourcery.com>
To: Jakub Jelinek <jakub@redhat.com>, <gcc-patches@gcc.gnu.org>
Subject: Re: C/C++ OpenACC: acc_pcopyin, acc_pcreate
Date: Wed, 24 May 2017 10:57:00 -0000	[thread overview]
Message-ID: <8737buy15s.fsf@hertz.schwinge.homeip.net> (raw)
In-Reply-To: <20170523110715.GQ8499@tucnak>

Hi Jakub!

On Tue, 23 May 2017 13:07:15 +0200, Jakub Jelinek <jakub@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 <thomas@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

  reply	other threads:[~2017-05-24 10:55 UTC|newest]

Thread overview: 8+ messages / expand[flat|nested]  mbox.gz  Atom feed  top
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: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 [this message]
2017-05-24 11:24     ` Jakub Jelinek
2017-05-24 13:41       ` Thomas Schwinge

Reply instructions:

You may reply publicly to this message via plain-text email
using any one of the following methods:

* Save the following mbox file, import it into your mail client,
  and reply-to-all from there: mbox

  Avoid top-posting and favor interleaved quoting:
  https://en.wikipedia.org/wiki/Posting_style#Interleaved_style

* Reply using the --to, --cc, and --in-reply-to
  switches of git-send-email(1):

  git send-email \
    --in-reply-to=8737buy15s.fsf@hertz.schwinge.homeip.net \
    --to=thomas@codesourcery.com \
    --cc=gcc-patches@gcc.gnu.org \
    --cc=jakub@redhat.com \
    /path/to/YOUR_REPLY

  https://kernel.org/pub/software/scm/git/docs/git-send-email.html

* If your mail client supports setting the In-Reply-To header
  via mailto: links, try the mailto: link
Be sure your reply has a Subject: header at the top and a blank line before the message body.
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).