[PATCH 0/7 libgomp,amdgcn] AMD GCN Offloading Support

classic Classic list List threaded Threaded
27 messages Options
12
Reply | Threaded
Open this post in threaded view
|

[PATCH 0/7 libgomp,amdgcn] AMD GCN Offloading Support

Andrew Stubbs-4
Hi all,

This patch series contributes initial OpenMP and OpenACC support for AMD
GCN GPUs.

The test results are not yet perfect, but there are many more passes than
failures, so this is a good starting point.  The rest of the issues can
be addressed as bugs during stage 3.

I have another, unfinished, patch to massage the testsuite itself.  I'll
post this shortly, once I've finished checking the forward port is
appropriate.

This series implements only single-worker support for OpenACC.  Julian
Brown may post the multiple-worker support soon, if it isn't too
difficult to forward-port.  Otherwise that will have to wait for GCC 11.

Andrew

Andrew Stubbs (7):
  Move generic libgomp files from nvptx to accel
  GCN mkoffload
  Add device number to GOMP_OFFLOAD_openacc_async_construct
  GCN libgomp port
  Optimize GCN OpenMP malloc performance
  Use a single worker for OpenACC on AMD GCN
  GCN Libgomp Plugin

 gcc/config/gcn/gcn.c                          |    4 +-
 gcc/config/gcn/gcn.opt                        |    2 +-
 gcc/config/gcn/mkoffload.c                    |  694 +++
 gcc/config/gcn/offload.h                      |   35 +
 include/gomp-constants.h                      |    2 +
 libgomp/Makefile.am                           |    2 +-
 libgomp/Makefile.in                           |   61 +-
 libgomp/config.h.in                           |    3 +
 .../config/{nvptx => accel}/libgomp-plugin.c  |    0
 libgomp/config/{nvptx => accel}/lock.c        |    0
 libgomp/config/{nvptx => accel}/mutex.c       |    0
 libgomp/config/{nvptx => accel}/mutex.h       |    0
 libgomp/config/{nvptx => accel}/oacc-async.c  |    0
 libgomp/config/{nvptx => accel}/oacc-cuda.c   |    0
 libgomp/config/{nvptx => accel}/oacc-host.c   |    0
 libgomp/config/{nvptx => accel}/oacc-init.c   |    0
 libgomp/config/{nvptx => accel}/oacc-mem.c    |    0
 libgomp/config/{nvptx => accel}/oacc-plugin.c |    0
 libgomp/config/{nvptx => accel}/omp-lock.h    |    0
 libgomp/config/{nvptx => accel}/openacc.f90   |    1 +
 libgomp/config/{nvptx => accel}/pool.h        |    0
 libgomp/config/{nvptx => accel}/proc.c        |    0
 libgomp/config/{nvptx => accel}/ptrlock.c     |    0
 libgomp/config/{nvptx => accel}/ptrlock.h     |    0
 libgomp/config/{nvptx => accel}/sem.c         |    0
 libgomp/config/{nvptx => accel}/sem.h         |    0
 .../{nvptx => accel}/thread-stacksize.h       |    0
 libgomp/config/gcn/affinity-fmt.c             |   51 +
 libgomp/config/gcn/bar.c                      |  232 +
 libgomp/config/gcn/bar.h                      |  168 +
 libgomp/config/gcn/doacross.h                 |   58 +
 libgomp/config/gcn/icv-device.c               |   72 +
 libgomp/config/gcn/oacc-target.c              |   31 +
 libgomp/config/gcn/simple-bar.h               |   61 +
 libgomp/config/gcn/target.c                   |   67 +
 libgomp/config/gcn/task.c                     |   39 +
 libgomp/config/gcn/team.c                     |  212 +
 libgomp/config/gcn/time.c                     |   52 +
 libgomp/configure                             |   61 +-
 libgomp/configure.ac                          |    2 +-
 libgomp/configure.tgt                         |    6 +-
 libgomp/libgomp-plugin.h                      |    5 +-
 libgomp/libgomp.h                             |   74 +
 libgomp/oacc-async.c                          |    3 +-
 libgomp/oacc-host.c                           |    2 +-
 libgomp/oacc-int.h                            |    9 +-
 libgomp/oacc-target.c                         |    1 +
 libgomp/openacc.f90                           |    1 +
 libgomp/openacc.h                             |    1 +
 libgomp/plugin/Makefrag.am                    |   14 +
 libgomp/plugin/configfrag.ac                  |   35 +
 libgomp/plugin/plugin-gcn.c                   | 3985 +++++++++++++++++
 libgomp/plugin/plugin-nvptx.c                 |    2 +-
 libgomp/team.c                                |   15 +-
 libgomp/work.c                                |    4 +-
 55 files changed, 6035 insertions(+), 32 deletions(-)
 create mode 100644 gcc/config/gcn/mkoffload.c
 create mode 100644 gcc/config/gcn/offload.h
 rename libgomp/config/{nvptx => accel}/libgomp-plugin.c (100%)
 rename libgomp/config/{nvptx => accel}/lock.c (100%)
 rename libgomp/config/{nvptx => accel}/mutex.c (100%)
 rename libgomp/config/{nvptx => accel}/mutex.h (100%)
 rename libgomp/config/{nvptx => accel}/oacc-async.c (100%)
 rename libgomp/config/{nvptx => accel}/oacc-cuda.c (100%)
 rename libgomp/config/{nvptx => accel}/oacc-host.c (100%)
 rename libgomp/config/{nvptx => accel}/oacc-init.c (100%)
 rename libgomp/config/{nvptx => accel}/oacc-mem.c (100%)
 rename libgomp/config/{nvptx => accel}/oacc-plugin.c (100%)
 rename libgomp/config/{nvptx => accel}/omp-lock.h (100%)
 rename libgomp/config/{nvptx => accel}/openacc.f90 (98%)
 rename libgomp/config/{nvptx => accel}/pool.h (100%)
 rename libgomp/config/{nvptx => accel}/proc.c (100%)
 rename libgomp/config/{nvptx => accel}/ptrlock.c (100%)
 rename libgomp/config/{nvptx => accel}/ptrlock.h (100%)
 rename libgomp/config/{nvptx => accel}/sem.c (100%)
 rename libgomp/config/{nvptx => accel}/sem.h (100%)
 rename libgomp/config/{nvptx => accel}/thread-stacksize.h (100%)
 create mode 100644 libgomp/config/gcn/affinity-fmt.c
 create mode 100644 libgomp/config/gcn/bar.c
 create mode 100644 libgomp/config/gcn/bar.h
 create mode 100644 libgomp/config/gcn/doacross.h
 create mode 100644 libgomp/config/gcn/icv-device.c
 create mode 100644 libgomp/config/gcn/oacc-target.c
 create mode 100644 libgomp/config/gcn/simple-bar.h
 create mode 100644 libgomp/config/gcn/target.c
 create mode 100644 libgomp/config/gcn/task.c
 create mode 100644 libgomp/config/gcn/team.c
 create mode 100644 libgomp/config/gcn/time.c
 create mode 100644 libgomp/oacc-target.c
 create mode 100644 libgomp/plugin/plugin-gcn.c

--
2.20.1

Reply | Threaded
Open this post in threaded view
|

[PATCH 1/7 libgomp,nvptx] Move generic libgomp files from nvptx to accel

Andrew Stubbs-4
This patch shouldn't change anything much at all; it's just an internal
reorganization of files.

The idea is to move files from libgomp "nvptx" directory that have
nothing NVPTX-specific in them.  By placing them in a separate "accel"
directory they can be shared with the GCN port, thus preventing much of
the duplication.

OK to commit?

Thanks

Andrew

2019-11-12  Andrew Stubbs  <[hidden email]>

        libgomp/
        * configure.tgt (nvptx*-*-*): Add "accel" directory.
        * config/nvptx/libgomp-plugin.c: Move ...
        * config/accel/libgomp-plugin.c: ... to here.
        * config/nvptx/lock.c: Move ...
        * config/accel/lock.c: ... to here.
        * config/nvptx/mutex.c: Move ...
        * config/accel/mutex.c: ... to here.
        * config/nvptx/mutex.h: Move ...
        * config/accel/mutex.h: ... to here.
        * config/nvptx/oacc-async.c: Move ...
        * config/accel/oacc-async.c: ... to here.
        * config/nvptx/oacc-cuda.c: Move ...
        * config/accel/oacc-cuda.c: ... to here.
        * config/nvptx/oacc-host.c: Move ...
        * config/accel/oacc-host.c: ... to here.
        * config/nvptx/oacc-init.c: Move ...
        * config/accel/oacc-init.c: ... to here.
        * config/nvptx/oacc-mem.c: Move ...
        * config/accel/oacc-mem.c: ... to here.
        * config/nvptx/oacc-plugin.c: Move ...
        * config/accel/oacc-plugin.c: ... to here.
        * config/nvptx/omp-lock.h: Move ...
        * config/accel/omp-lock.h: ... to here.
        * config/nvptx/openacc.f90: Move ...
        * config/accel/openacc.f90: ... to here.
        * config/nvptx/pool.h: Move ...
        * config/accel/pool.h: ... to here.
        * config/nvptx/proc.c: Move ...
        * config/accel/proc.c: ... to here.
        * config/nvptx/ptrlock.c: Move ...
        * config/accel/ptrlock.c: ... to here.
        * config/nvptx/ptrlock.h: Move ...
        * config/accel/ptrlock.h: ... to here.
        * config/nvptx/sem.c: Move ...
        * config/accel/sem.c: ... to here.
        * config/nvptx/sem.h: Move ...
        * config/accel/sem.h: ... to here.
        * config/nvptx/thread-stacksize.h: Move ...
        * config/accel/thread-stacksize.h: ... to here.
---
 libgomp/config/{nvptx => accel}/libgomp-plugin.c   | 0
 libgomp/config/{nvptx => accel}/lock.c             | 0
 libgomp/config/{nvptx => accel}/mutex.c            | 0
 libgomp/config/{nvptx => accel}/mutex.h            | 0
 libgomp/config/{nvptx => accel}/oacc-async.c       | 0
 libgomp/config/{nvptx => accel}/oacc-cuda.c        | 0
 libgomp/config/{nvptx => accel}/oacc-host.c        | 0
 libgomp/config/{nvptx => accel}/oacc-init.c        | 0
 libgomp/config/{nvptx => accel}/oacc-mem.c         | 0
 libgomp/config/{nvptx => accel}/oacc-plugin.c      | 0
 libgomp/config/{nvptx => accel}/omp-lock.h         | 0
 libgomp/config/{nvptx => accel}/openacc.f90        | 0
 libgomp/config/{nvptx => accel}/pool.h             | 0
 libgomp/config/{nvptx => accel}/proc.c             | 0
 libgomp/config/{nvptx => accel}/ptrlock.c          | 0
 libgomp/config/{nvptx => accel}/ptrlock.h          | 0
 libgomp/config/{nvptx => accel}/sem.c              | 0
 libgomp/config/{nvptx => accel}/sem.h              | 0
 libgomp/config/{nvptx => accel}/thread-stacksize.h | 0
 libgomp/configure.tgt                              | 2 +-
 20 files changed, 1 insertion(+), 1 deletion(-)
 rename libgomp/config/{nvptx => accel}/libgomp-plugin.c (100%)
 rename libgomp/config/{nvptx => accel}/lock.c (100%)
 rename libgomp/config/{nvptx => accel}/mutex.c (100%)
 rename libgomp/config/{nvptx => accel}/mutex.h (100%)
 rename libgomp/config/{nvptx => accel}/oacc-async.c (100%)
 rename libgomp/config/{nvptx => accel}/oacc-cuda.c (100%)
 rename libgomp/config/{nvptx => accel}/oacc-host.c (100%)
 rename libgomp/config/{nvptx => accel}/oacc-init.c (100%)
 rename libgomp/config/{nvptx => accel}/oacc-mem.c (100%)
 rename libgomp/config/{nvptx => accel}/oacc-plugin.c (100%)
 rename libgomp/config/{nvptx => accel}/omp-lock.h (100%)
 rename libgomp/config/{nvptx => accel}/openacc.f90 (100%)
 rename libgomp/config/{nvptx => accel}/pool.h (100%)
 rename libgomp/config/{nvptx => accel}/proc.c (100%)
 rename libgomp/config/{nvptx => accel}/ptrlock.c (100%)
 rename libgomp/config/{nvptx => accel}/ptrlock.h (100%)
 rename libgomp/config/{nvptx => accel}/sem.c (100%)
 rename libgomp/config/{nvptx => accel}/sem.h (100%)
 rename libgomp/config/{nvptx => accel}/thread-stacksize.h (100%)


diff --git a/libgomp/config/nvptx/libgomp-plugin.c b/libgomp/config/accel/libgomp-plugin.c
similarity index 100%
rename from libgomp/config/nvptx/libgomp-plugin.c
rename to libgomp/config/accel/libgomp-plugin.c
diff --git a/libgomp/config/nvptx/lock.c b/libgomp/config/accel/lock.c
similarity index 100%
rename from libgomp/config/nvptx/lock.c
rename to libgomp/config/accel/lock.c
diff --git a/libgomp/config/nvptx/mutex.c b/libgomp/config/accel/mutex.c
similarity index 100%
rename from libgomp/config/nvptx/mutex.c
rename to libgomp/config/accel/mutex.c
diff --git a/libgomp/config/nvptx/mutex.h b/libgomp/config/accel/mutex.h
similarity index 100%
rename from libgomp/config/nvptx/mutex.h
rename to libgomp/config/accel/mutex.h
diff --git a/libgomp/config/nvptx/oacc-async.c b/libgomp/config/accel/oacc-async.c
similarity index 100%
rename from libgomp/config/nvptx/oacc-async.c
rename to libgomp/config/accel/oacc-async.c
diff --git a/libgomp/config/nvptx/oacc-cuda.c b/libgomp/config/accel/oacc-cuda.c
similarity index 100%
rename from libgomp/config/nvptx/oacc-cuda.c
rename to libgomp/config/accel/oacc-cuda.c
diff --git a/libgomp/config/nvptx/oacc-host.c b/libgomp/config/accel/oacc-host.c
similarity index 100%
rename from libgomp/config/nvptx/oacc-host.c
rename to libgomp/config/accel/oacc-host.c
diff --git a/libgomp/config/nvptx/oacc-init.c b/libgomp/config/accel/oacc-init.c
similarity index 100%
rename from libgomp/config/nvptx/oacc-init.c
rename to libgomp/config/accel/oacc-init.c
diff --git a/libgomp/config/nvptx/oacc-mem.c b/libgomp/config/accel/oacc-mem.c
similarity index 100%
rename from libgomp/config/nvptx/oacc-mem.c
rename to libgomp/config/accel/oacc-mem.c
diff --git a/libgomp/config/nvptx/oacc-plugin.c b/libgomp/config/accel/oacc-plugin.c
similarity index 100%
rename from libgomp/config/nvptx/oacc-plugin.c
rename to libgomp/config/accel/oacc-plugin.c
diff --git a/libgomp/config/nvptx/omp-lock.h b/libgomp/config/accel/omp-lock.h
similarity index 100%
rename from libgomp/config/nvptx/omp-lock.h
rename to libgomp/config/accel/omp-lock.h
diff --git a/libgomp/config/nvptx/openacc.f90 b/libgomp/config/accel/openacc.f90
similarity index 100%
rename from libgomp/config/nvptx/openacc.f90
rename to libgomp/config/accel/openacc.f90
diff --git a/libgomp/config/nvptx/pool.h b/libgomp/config/accel/pool.h
similarity index 100%
rename from libgomp/config/nvptx/pool.h
rename to libgomp/config/accel/pool.h
diff --git a/libgomp/config/nvptx/proc.c b/libgomp/config/accel/proc.c
similarity index 100%
rename from libgomp/config/nvptx/proc.c
rename to libgomp/config/accel/proc.c
diff --git a/libgomp/config/nvptx/ptrlock.c b/libgomp/config/accel/ptrlock.c
similarity index 100%
rename from libgomp/config/nvptx/ptrlock.c
rename to libgomp/config/accel/ptrlock.c
diff --git a/libgomp/config/nvptx/ptrlock.h b/libgomp/config/accel/ptrlock.h
similarity index 100%
rename from libgomp/config/nvptx/ptrlock.h
rename to libgomp/config/accel/ptrlock.h
diff --git a/libgomp/config/nvptx/sem.c b/libgomp/config/accel/sem.c
similarity index 100%
rename from libgomp/config/nvptx/sem.c
rename to libgomp/config/accel/sem.c
diff --git a/libgomp/config/nvptx/sem.h b/libgomp/config/accel/sem.h
similarity index 100%
rename from libgomp/config/nvptx/sem.h
rename to libgomp/config/accel/sem.h
diff --git a/libgomp/config/nvptx/thread-stacksize.h b/libgomp/config/accel/thread-stacksize.h
similarity index 100%
rename from libgomp/config/nvptx/thread-stacksize.h
rename to libgomp/config/accel/thread-stacksize.h
diff --git a/libgomp/configure.tgt b/libgomp/configure.tgt
index b88bf72fe3d..c5ae9a9e39a 100644
--- a/libgomp/configure.tgt
+++ b/libgomp/configure.tgt
@@ -154,7 +154,7 @@ case "${target}" in
  ;;
 
   nvptx*-*-*)
- config_path="nvptx"
+ config_path="nvptx accel"
  ;;
 
   *-*-rtems*)
Reply | Threaded
Open this post in threaded view
|

[PATCH 2/7 amdgcn] GCN mkoffload

Andrew Stubbs-4
In reply to this post by Andrew Stubbs-4
This patch adds the mkoffload tool to the amdgcn backend.  It's similar,
but not quite the same as that on the openacc-gcc-9-branch.

I will commit this patch when the others in this series are approved.

Andrew


2019-11-12  Andrew Stubbs  <[hidden email]>

        gcc/
        * config/gcn/mkoffload.c: New file.
        * config/gcn/offload.h: New file.
---
 gcc/config/gcn/mkoffload.c | 694 +++++++++++++++++++++++++++++++++++++
 gcc/config/gcn/offload.h   |  35 ++
 2 files changed, 729 insertions(+)
 create mode 100644 gcc/config/gcn/mkoffload.c
 create mode 100644 gcc/config/gcn/offload.h


diff --git a/gcc/config/gcn/mkoffload.c b/gcc/config/gcn/mkoffload.c
new file mode 100644
index 00000000000..40b56375b75
--- /dev/null
+++ b/gcc/config/gcn/mkoffload.c
@@ -0,0 +1,694 @@
+/* Offload image generation tool for AMD GCN.
+
+   Copyright (C) 2014-2019 Free Software Foundation, Inc.
+
+   This file is part of GCC.
+
+   GCC is free software; you can redistribute it and/or modify it
+   under the terms of the GNU General Public License as published
+   by the Free Software Foundation; either version 3, or (at your
+   option) any later version.
+
+   GCC is distributed in the hope that it will be useful, but WITHOUT
+   ANY WARRANTY; without even the implied warranty of MERCHANTABILITY
+   or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU General Public
+   License for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with GCC; see the file COPYING3.  If not see
+   <http://www.gnu.org/licenses/>.  */
+
+/* Munges GCN assembly into a C source file defining the GCN code as a
+   string.
+
+   This is not a complete assembler.  We presume the source is well
+   formed from the compiler and can die horribly if it is not.  */
+
+#include "config.h"
+#include "system.h"
+#include "coretypes.h"
+#include "obstack.h"
+#include "diagnostic.h"
+#include "intl.h"
+#include <libgen.h>
+#include "collect-utils.h"
+#include "gomp-constants.h"
+
+const char tool_name[] = "gcn mkoffload";
+
+#define COMMENT_PREFIX "#"
+
+struct id_map
+{
+  id_map *next;
+  char *gcn_name;
+};
+
+static id_map *func_ids, **funcs_tail = &func_ids;
+static id_map *var_ids, **vars_tail = &var_ids;
+
+/* Files to unlink.  */
+static const char *gcn_s1_name;
+static const char *gcn_s2_name;
+static const char *gcn_o_name;
+static const char *gcn_cfile_name;
+
+enum offload_abi offload_abi = OFFLOAD_ABI_UNSET;
+
+/* Delete tempfiles.  */
+
+void
+tool_cleanup (bool from_signal ATTRIBUTE_UNUSED)
+{
+  if (gcn_cfile_name)
+    maybe_unlink (gcn_cfile_name);
+  if (gcn_s1_name)
+    maybe_unlink (gcn_s1_name);
+  if (gcn_s2_name)
+    maybe_unlink (gcn_s2_name);
+  if (gcn_o_name)
+    maybe_unlink (gcn_o_name);
+}
+
+static void
+mkoffload_cleanup (void)
+{
+  tool_cleanup (false);
+}
+
+/* Unlink FILE unless requested otherwise.  */
+
+void
+maybe_unlink (const char *file)
+{
+  if (!save_temps)
+    {
+      if (unlink_if_ordinary (file) && errno != ENOENT)
+ fatal_error (input_location, "deleting file %s: %m", file);
+    }
+  else if (verbose)
+    fprintf (stderr, "[Leaving %s]\n", file);
+}
+
+/* Add or change the value of an environment variable, outputting the
+   change to standard error if in verbose mode.  */
+
+static void
+xputenv (const char *string)
+{
+  if (verbose)
+    fprintf (stderr, "%s\n", string);
+  putenv (CONST_CAST (char *, string));
+}
+
+/* Read the whole input file.  It will be NUL terminated (but
+   remember, there could be a NUL in the file itself.  */
+
+static const char *
+read_file (FILE *stream, size_t *plen)
+{
+  size_t alloc = 16384;
+  size_t base = 0;
+  char *buffer;
+
+  if (!fseek (stream, 0, SEEK_END))
+    {
+      /* Get the file size.  */
+      long s = ftell (stream);
+      if (s >= 0)
+ alloc = s + 100;
+      fseek (stream, 0, SEEK_SET);
+    }
+  buffer = XNEWVEC (char, alloc);
+
+  for (;;)
+    {
+      size_t n = fread (buffer + base, 1, alloc - base - 1, stream);
+
+      if (!n)
+ break;
+      base += n;
+      if (base + 1 == alloc)
+ {
+  alloc *= 2;
+  buffer = XRESIZEVEC (char, buffer, alloc);
+ }
+    }
+  buffer[base] = 0;
+  *plen = base;
+  return buffer;
+}
+
+/* Parse STR, saving found tokens into PVALUES and return their number.
+   Tokens are assumed to be delimited by ':'.  */
+
+static unsigned
+parse_env_var (const char *str, char ***pvalues)
+{
+  const char *curval, *nextval;
+  char **values;
+  unsigned num = 1, i;
+
+  curval = strchr (str, ':');
+  while (curval)
+    {
+      num++;
+      curval = strchr (curval + 1, ':');
+    }
+
+  values = (char **) xmalloc (num * sizeof (char *));
+  curval = str;
+  nextval = strchr (curval, ':');
+  if (nextval == NULL)
+    nextval = strchr (curval, '\0');
+
+  for (i = 0; i < num; i++)
+    {
+      int l = nextval - curval;
+      values[i] = (char *) xmalloc (l + 1);
+      memcpy (values[i], curval, l);
+      values[i][l] = 0;
+      curval = nextval + 1;
+      nextval = strchr (curval, ':');
+      if (nextval == NULL)
+ nextval = strchr (curval, '\0');
+    }
+  *pvalues = values;
+  return num;
+}
+
+/* Auxiliary function that frees elements of PTR and PTR itself.
+   N is number of elements to be freed.  If PTR is NULL, nothing is freed.
+   If an element is NULL, subsequent elements are not freed.  */
+
+static void
+free_array_of_ptrs (void **ptr, unsigned n)
+{
+  unsigned i;
+  if (!ptr)
+    return;
+  for (i = 0; i < n; i++)
+    {
+      if (!ptr[i])
+ break;
+      free (ptr[i]);
+    }
+  free (ptr);
+  return;
+}
+
+/* Check whether NAME can be accessed in MODE.  This is like access,
+   except that it never considers directories to be executable.  */
+
+static int
+access_check (const char *name, int mode)
+{
+  if (mode == X_OK)
+    {
+      struct stat st;
+
+      if (stat (name, &st) < 0 || S_ISDIR (st.st_mode))
+ return -1;
+    }
+
+  return access (name, mode);
+}
+
+/* Parse an input assembler file, extract the offload tables etc.,
+   and output (1) the assembler code, minus the tables (which can contain
+   problematic relocations), and (2) a C file with the offload tables
+   encoded as structured data.  */
+
+static void
+process_asm (FILE *in, FILE *out, FILE *cfile)
+{
+  int fn_count = 0, var_count = 0, dims_count = 0;
+  struct obstack fns_os, vars_os, varsizes_os, dims_os;
+  obstack_init (&fns_os);
+  obstack_init (&vars_os);
+  obstack_init (&varsizes_os);
+  obstack_init (&dims_os);
+
+  struct oaccdims
+  {
+    int d[3];
+    char *name;
+  } dim;
+
+  /* Always add _init_array and _fini_array as kernels.  */
+  obstack_ptr_grow (&fns_os, xstrdup ("_init_array"));
+  obstack_ptr_grow (&fns_os, xstrdup ("_fini_array"));
+  fn_count += 2;
+
+  char buf[1000];
+  enum { IN_CODE, IN_VARS, IN_FUNCS } state = IN_CODE;
+  while (fgets (buf, sizeof (buf), in))
+    {
+      switch (state)
+ {
+ case IN_CODE:
+  {
+    if (sscanf (buf, " ;; OPENACC-DIMS: %d, %d, %d : %ms\n",
+ &dim.d[0], &dim.d[1], &dim.d[2], &dim.name) == 4)
+      {
+ obstack_grow (&dims_os, &dim, sizeof (dim));
+ dims_count++;
+      }
+    break;
+  }
+ case IN_VARS:
+  {
+    char *varname;
+    unsigned varsize;
+    if (sscanf (buf, " .8byte %ms\n", &varname))
+      {
+ obstack_ptr_grow (&vars_os, varname);
+ fgets (buf, sizeof (buf), in);
+ if (!sscanf (buf, " .8byte %u\n", &varsize))
+  abort ();
+ obstack_int_grow (&varsizes_os, varsize);
+ var_count++;
+
+ /* The HSA Runtime cannot locate the symbol if it is not
+   exported from the kernel.  */
+ fprintf (out, "\t.global %s\n", varname);
+      }
+    break;
+  }
+ case IN_FUNCS:
+  {
+    char *funcname;
+    if (sscanf (buf, "\t.8byte\t%ms\n", &funcname))
+      {
+ obstack_ptr_grow (&fns_os, funcname);
+ fn_count++;
+ continue;
+      }
+    break;
+  }
+ }
+
+      char dummy;
+      if (sscanf (buf, " .section .gnu.offload_vars%c", &dummy) > 0)
+ state = IN_VARS;
+      else if (sscanf (buf, " .section .gnu.offload_funcs%c", &dummy) > 0)
+ state = IN_FUNCS;
+      else if (sscanf (buf, " .section %c", &dummy) > 0
+       || sscanf (buf, " .text%c", &dummy) > 0
+       || sscanf (buf, " .bss%c", &dummy) > 0
+       || sscanf (buf, " .data%c", &dummy) > 0
+       || sscanf (buf, " .ident %c", &dummy) > 0)
+ state = IN_CODE;
+
+      if (state == IN_CODE)
+ fputs (buf, out);
+    }
+
+  char **fns = XOBFINISH (&fns_os, char **);
+  struct oaccdims *dims = XOBFINISH (&dims_os, struct oaccdims *);
+
+  fprintf (cfile, "#include <stdlib.h>\n");
+  fprintf (cfile, "#include <stdbool.h>\n\n");
+
+  char **vars = XOBFINISH (&vars_os, char **);
+  unsigned *varsizes = XOBFINISH (&varsizes_os, unsigned *);
+  fprintf (cfile,
+   "static const struct global_var_info {\n"
+   "  const char *name;\n"
+   "  void *address;\n"
+   "} vars[] = {\n");
+  int i;
+  for (i = 0; i < var_count; ++i)
+    {
+      const char *sep = i < var_count - 1 ? "," : " ";
+      fprintf (cfile, "  { \"%s\", NULL }%s /* size: %u */\n", vars[i], sep,
+       varsizes[i]);
+    }
+  fprintf (cfile, "};\n\n");
+
+  obstack_free (&vars_os, NULL);
+  obstack_free (&varsizes_os, NULL);
+
+  /* Dump out function idents.  */
+  fprintf (cfile, "static const struct hsa_kernel_description {\n"
+   "  const char *name;\n"
+   "  int oacc_dims[3];\n"
+   "} gcn_kernels[] = {\n  ");
+  dim.d[0] = dim.d[1] = dim.d[2] = 0;
+  const char *comma;
+  for (comma = "", i = 0; i < fn_count; comma = ",\n  ", i++)
+    {
+      /* Find if we recorded dimensions for this function.  */
+      int *d = dim.d; /* Previously zeroed.  */
+      for (int j = 0; j < dims_count; j++)
+ if (strcmp (fns[i], dims[j].name) == 0)
+  {
+    d = dims[j].d;
+    break;
+  }
+
+      fprintf (cfile, "%s{\"%s\", {%d, %d, %d}}", comma,
+       fns[i], d[0], d[1], d[2]);
+
+      free (fns[i]);
+    }
+  fprintf (cfile, "\n};\n\n");
+
+  obstack_free (&fns_os, NULL);
+  for (i = 0; i < dims_count; i++)
+    free (dims[i].name);
+  obstack_free (&dims_os, NULL);
+}
+
+/* Embed an object file into a C source file.  */
+
+static void
+process_obj (FILE *in, FILE *cfile)
+{
+  size_t len = 0;
+  const char *input = read_file (in, &len);
+
+  /* Dump out an array containing the binary.
+     FIXME: do this with objcopy.  */
+  fprintf (cfile, "static unsigned char gcn_code[] = {");
+  for (size_t i = 0; i < len; i += 17)
+    {
+      fprintf (cfile, "\n\t");
+      for (size_t j = i; j < i + 17 && j < len; j++)
+ fprintf (cfile, "%3u,", (unsigned char) input[j]);
+    }
+  fprintf (cfile, "\n};\n\n");
+
+  fprintf (cfile,
+   "static const struct gcn_image {\n"
+   "  size_t size;\n"
+   "  void *image;\n"
+   "} gcn_image = {\n"
+   "  %zu,\n"
+   "  gcn_code\n"
+   "};\n\n",
+   len);
+
+  fprintf (cfile,
+   "static const struct gcn_image_desc {\n"
+   "  const struct gcn_image *gcn_image;\n"
+   "  unsigned kernel_count;\n"
+   "  const struct hsa_kernel_description *kernel_infos;\n"
+   "  unsigned global_variable_count;\n"
+   "  const struct global_var_info *global_variables;\n"
+   "} target_data = {\n"
+   "  &gcn_image,\n"
+   "  sizeof (gcn_kernels) / sizeof (gcn_kernels[0]),\n"
+   "  gcn_kernels,\n"
+   "  sizeof (vars) / sizeof (vars[0]),\n"
+   "  vars\n"
+   "};\n\n");
+
+  fprintf (cfile,
+   "#ifdef __cplusplus\n"
+   "extern \"C\" {\n"
+   "#endif\n"
+   "extern void GOMP_offload_register_ver"
+   " (unsigned, const void *, int, const void *);\n"
+   "extern void GOMP_offload_unregister_ver"
+   " (unsigned, const void *, int, const void *);\n"
+   "#ifdef __cplusplus\n"
+   "}\n"
+   "#endif\n\n");
+
+  fprintf (cfile, "extern const void *const __OFFLOAD_TABLE__[];\n\n");
+
+  fprintf (cfile, "static __attribute__((constructor)) void init (void)\n"
+   "{\n"
+   "  GOMP_offload_register_ver (%#x, __OFFLOAD_TABLE__,"
+   " %d/*GCN*/, &target_data);\n"
+   "};\n",
+   GOMP_VERSION_PACK (GOMP_VERSION, GOMP_VERSION_GCN),
+   GOMP_DEVICE_GCN);
+
+  fprintf (cfile, "static __attribute__((destructor)) void fini (void)\n"
+   "{\n"
+   "  GOMP_offload_unregister_ver (%#x, __OFFLOAD_TABLE__,"
+   " %d/*GCN*/, &target_data);\n"
+   "};\n",
+   GOMP_VERSION_PACK (GOMP_VERSION, GOMP_VERSION_GCN),
+   GOMP_DEVICE_GCN);
+}
+
+/* Compile a C file using the host compiler.  */
+
+static void
+compile_native (const char *infile, const char *outfile, const char *compiler)
+{
+  const char *collect_gcc_options = getenv ("COLLECT_GCC_OPTIONS");
+  if (!collect_gcc_options)
+    fatal_error (input_location,
+ "environment variable COLLECT_GCC_OPTIONS must be set");
+
+  struct obstack argv_obstack;
+  obstack_init (&argv_obstack);
+  obstack_ptr_grow (&argv_obstack, compiler);
+  if (save_temps)
+    obstack_ptr_grow (&argv_obstack, "-save-temps");
+  if (verbose)
+    obstack_ptr_grow (&argv_obstack, "-v");
+  switch (offload_abi)
+    {
+    case OFFLOAD_ABI_LP64:
+      obstack_ptr_grow (&argv_obstack, "-m64");
+      break;
+    case OFFLOAD_ABI_ILP32:
+      obstack_ptr_grow (&argv_obstack, "-m32");
+      break;
+    default:
+      gcc_unreachable ();
+    }
+  obstack_ptr_grow (&argv_obstack, infile);
+  obstack_ptr_grow (&argv_obstack, "-c");
+  obstack_ptr_grow (&argv_obstack, "-o");
+  obstack_ptr_grow (&argv_obstack, outfile);
+  obstack_ptr_grow (&argv_obstack, NULL);
+
+  const char **new_argv = XOBFINISH (&argv_obstack, const char **);
+  fork_execute (new_argv[0], CONST_CAST (char **, new_argv), true);
+  obstack_free (&argv_obstack, NULL);
+}
+
+int
+main (int argc, char **argv)
+{
+  FILE *in = stdin;
+  FILE *out = stdout;
+  FILE *cfile = stdout;
+  const char *outname = 0, *offloadsrc = 0;
+
+  progname = "mkoffload";
+  diagnostic_initialize (global_dc, 0);
+
+  if (atexit (mkoffload_cleanup) != 0)
+    fatal_error (input_location, "atexit failed");
+
+  char *collect_gcc = getenv ("COLLECT_GCC");
+  if (collect_gcc == NULL)
+    fatal_error (input_location, "COLLECT_GCC must be set.");
+  const char *gcc_path = dirname (ASTRDUP (collect_gcc));
+  const char *gcc_exec = basename (ASTRDUP (collect_gcc));
+
+  size_t len = (strlen (gcc_path) + 1 + strlen (GCC_INSTALL_NAME) + 1);
+  char *driver = XALLOCAVEC (char, len);
+
+  if (strcmp (gcc_exec, collect_gcc) == 0)
+    /* collect_gcc has no path, so it was found in PATH.  Make sure we also
+       find accel-gcc in PATH.  */
+    gcc_path = NULL;
+
+  int driver_used = 0;
+  if (gcc_path != NULL)
+    driver_used = sprintf (driver, "%s/", gcc_path);
+  sprintf (driver + driver_used, "%s", GCC_INSTALL_NAME);
+
+  bool found = false;
+  if (gcc_path == NULL)
+    found = true;
+  else if (access_check (driver, X_OK) == 0)
+    found = true;
+  else
+    {
+      /* Don't use alloca pointer with XRESIZEVEC.  */
+      driver = NULL;
+      /* Look in all COMPILER_PATHs for GCC_INSTALL_NAME.  */
+      char **paths = NULL;
+      unsigned n_paths;
+      n_paths = parse_env_var (getenv ("COMPILER_PATH"), &paths);
+      for (unsigned i = 0; i < n_paths; i++)
+ {
+  len = strlen (paths[i]) + 1 + strlen (GCC_INSTALL_NAME) + 1;
+  driver = XRESIZEVEC (char, driver, len);
+  sprintf (driver, "%s/%s", paths[i], GCC_INSTALL_NAME);
+  if (access_check (driver, X_OK) == 0)
+    {
+      found = true;
+      break;
+    }
+ }
+      free_array_of_ptrs ((void **) paths, n_paths);
+    }
+
+  if (!found)
+    fatal_error (input_location,
+ "offload compiler %s not found", GCC_INSTALL_NAME);
+
+  /* We may be called with all the arguments stored in some file and
+     passed with @file.  Expand them into argv before processing.  */
+  expandargv (&argc, &argv);
+
+  /* Scan the argument vector.  */
+  bool fopenmp = false;
+  bool fopenacc = false;
+  for (int i = 1; i < argc; i++)
+    {
+#define STR "-foffload-abi="
+      if (strncmp (argv[i], STR, strlen (STR)) == 0)
+ {
+  if (strcmp (argv[i] + strlen (STR), "lp64") == 0)
+    offload_abi = OFFLOAD_ABI_LP64;
+  else if (strcmp (argv[i] + strlen (STR), "ilp32") == 0)
+    offload_abi = OFFLOAD_ABI_ILP32;
+  else
+    fatal_error (input_location,
+ "unrecognizable argument of option " STR);
+ }
+#undef STR
+      else if (strcmp (argv[i], "-fopenmp") == 0)
+ fopenmp = true;
+      else if (strcmp (argv[i], "-fopenacc") == 0)
+ fopenacc = true;
+      else if (strcmp (argv[i], "-save-temps") == 0)
+ save_temps = true;
+      else if (strcmp (argv[i], "-v") == 0)
+ verbose = true;
+    }
+  if (!(fopenacc ^ fopenmp))
+    fatal_error (input_location, "either -fopenacc or -fopenmp must be set");
+
+  const char *abi;
+  switch (offload_abi)
+    {
+    case OFFLOAD_ABI_LP64:
+      abi = "-m64";
+      break;
+    case OFFLOAD_ABI_ILP32:
+      abi = "-m32";
+      break;
+    default:
+      gcc_unreachable ();
+    }
+
+  gcn_s1_name = make_temp_file (".mkoffload.1.s");
+  gcn_s2_name = make_temp_file (".mkoffload.2.s");
+  gcn_o_name = make_temp_file (".mkoffload.hsaco");
+  gcn_cfile_name = make_temp_file (".c");
+
+  /* Build arguments for compiler pass.  */
+  struct obstack cc_argv_obstack;
+  obstack_init (&cc_argv_obstack);
+  obstack_ptr_grow (&cc_argv_obstack, driver);
+  obstack_ptr_grow (&cc_argv_obstack, "-S");
+
+  if (save_temps)
+    obstack_ptr_grow (&cc_argv_obstack, "-save-temps");
+  if (verbose)
+    obstack_ptr_grow (&cc_argv_obstack, "-v");
+  obstack_ptr_grow (&cc_argv_obstack, abi);
+  obstack_ptr_grow (&cc_argv_obstack, "-xlto");
+  if (fopenmp)
+    obstack_ptr_grow (&cc_argv_obstack, "-mgomp");
+
+  for (int ix = 1; ix != argc; ix++)
+    {
+      if (!strcmp (argv[ix], "-o") && ix + 1 != argc)
+ outname = argv[++ix];
+      else
+ {
+  obstack_ptr_grow (&cc_argv_obstack, argv[ix]);
+
+  if (argv[ix][0] != '-')
+    offloadsrc = argv[ix];
+ }
+    }
+
+  obstack_ptr_grow (&cc_argv_obstack, "-o");
+  obstack_ptr_grow (&cc_argv_obstack, gcn_s1_name);
+  obstack_ptr_grow (&cc_argv_obstack,
+    concat ("-mlocal-symbol-id=", offloadsrc, NULL));
+  obstack_ptr_grow (&cc_argv_obstack, NULL);
+  const char **cc_argv = XOBFINISH (&cc_argv_obstack, const char **);
+
+  /* Build arguments for assemble/link pass.  */
+  struct obstack ld_argv_obstack;
+  obstack_init (&ld_argv_obstack);
+  obstack_ptr_grow (&ld_argv_obstack, driver);
+  obstack_ptr_grow (&ld_argv_obstack, gcn_s2_name);
+  obstack_ptr_grow (&ld_argv_obstack, "-lgomp");
+
+  for (int i = 1; i < argc; i++)
+    if (strncmp (argv[i], "-l", 2) == 0
+ || strncmp (argv[i], "-Wl", 3) == 0
+ || strncmp (argv[i], "-march", 6) == 0)
+      obstack_ptr_grow (&ld_argv_obstack, argv[i]);
+
+  obstack_ptr_grow (&ld_argv_obstack, "-o");
+  obstack_ptr_grow (&ld_argv_obstack, gcn_o_name);
+  obstack_ptr_grow (&ld_argv_obstack, NULL);
+  const char **ld_argv = XOBFINISH (&ld_argv_obstack, const char **);
+
+  /* Clean up unhelpful environment variables.  */
+  char *execpath = getenv ("GCC_EXEC_PREFIX");
+  char *cpath = getenv ("COMPILER_PATH");
+  char *lpath = getenv ("LIBRARY_PATH");
+  unsetenv ("GCC_EXEC_PREFIX");
+  unsetenv ("COMPILER_PATH");
+  unsetenv ("LIBRARY_PATH");
+
+  /* Run the compiler pass.  */
+  fork_execute (cc_argv[0], CONST_CAST (char **, cc_argv), true);
+  obstack_free (&cc_argv_obstack, NULL);
+
+  in = fopen (gcn_s1_name, "r");
+  if (!in)
+    fatal_error (input_location, "cannot open intermediate gcn asm file");
+
+  out = fopen (gcn_s2_name, "w");
+  if (!out)
+    fatal_error (input_location, "cannot open '%s'", gcn_s2_name);
+
+  cfile = fopen (gcn_cfile_name, "w");
+  if (!cfile)
+    fatal_error (input_location, "cannot open '%s'", gcn_cfile_name);
+
+  process_asm (in, out, cfile);
+
+  fclose (in);
+  fclose (out);
+
+  /* Run the assemble/link pass.  */
+  fork_execute (ld_argv[0], CONST_CAST (char **, ld_argv), true);
+  obstack_free (&ld_argv_obstack, NULL);
+
+  in = fopen (gcn_o_name, "r");
+  if (!in)
+    fatal_error (input_location, "cannot open intermediate gcn obj file");
+
+  process_obj (in, cfile);
+
+  fclose (in);
+  fclose (cfile);
+
+  xputenv (concat ("GCC_EXEC_PREFIX=", execpath, NULL));
+  xputenv (concat ("COMPILER_PATH=", cpath, NULL));
+  xputenv (concat ("LIBRARY_PATH=", lpath, NULL));
+
+  compile_native (gcn_cfile_name, outname, collect_gcc);
+
+  return 0;
+}
diff --git a/gcc/config/gcn/offload.h b/gcc/config/gcn/offload.h
new file mode 100644
index 00000000000..795ee3f0e0e
--- /dev/null
+++ b/gcc/config/gcn/offload.h
@@ -0,0 +1,35 @@
+/* Support for AMD GCN offloading.
+
+   Copyright (C) 2014-2019 Free Software Foundation, Inc.
+
+   This file is part of GCC.
+
+   GCC is free software; you can redistribute it and/or modify
+   it under the terms of the GNU General Public License as published by
+   the Free Software Foundation; either version 3, or (at your option)
+   any later version.
+
+   GCC is distributed in the hope that it will be useful,
+   but WITHOUT ANY WARRANTY; without even the implied warranty of
+   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+   GNU General Public License for more details.
+
+   Under Section 7 of GPL version 3, you are granted additional
+   permissions described in the GCC Runtime Library Exception, version
+   3.1, as published by the Free Software Foundation.
+
+   You should have received a copy of the GNU General Public License and
+   a copy of the GCC Runtime Library Exception along with this program;
+   see the files COPYING3 and COPYING.RUNTIME respectively.  If not, see
+   <http://www.gnu.org/licenses/>.  */
+
+#ifndef GCC_GCN_OFFLOAD_H
+#define GCC_GCN_OFFLOAD_H
+
+/* Support for OpenACC acc_on_device.  */
+
+#include "gomp-constants.h"
+
+#define ACCEL_COMPILER_acc_device GOMP_DEVICE_GCN
+
+#endif
Reply | Threaded
Open this post in threaded view
|

[PATCH 3/7 libgomp,nvptx] Add device number to GOMP_OFFLOAD_openacc_async_construct

Andrew Stubbs-4
In reply to this post by Andrew Stubbs-4
This patch is preparatory for the amdgcn plugin.  The current
implementation was written for CUDA in which the device associated with
the queue is inferred by some hidden magic (which seems questionable to
me but then I don't fully understand it).  The GCN plugin needs to know
for which device the queue is intended, so this simply provides that
information to the queue constructor.

OK to commit?

Thanks

Andrew


2019-11-12  Andrew Stubbs  <[hidden email]>

        libgomp/
        * libgomp-plugin.h (GOMP_OFFLOAD_openacc_async_construct): Add int
        parameter.
        * oacc-async.c (lookup_goacc_asyncqueue): Pass device number to the
        queue constructor.
        * oacc-host.c (host_openacc_async_construct): Add device parameter.
        * plugin/plugin-nvptx.c (GOMP_OFFLOAD_openacc_async_construct): Add
        device parameter.
---
 libgomp/libgomp-plugin.h      | 2 +-
 libgomp/oacc-async.c          | 3 ++-
 libgomp/oacc-host.c           | 2 +-
 libgomp/plugin/plugin-nvptx.c | 2 +-
 4 files changed, 5 insertions(+), 4 deletions(-)


diff --git a/libgomp/libgomp-plugin.h b/libgomp/libgomp-plugin.h
index 01483f27f4c..de969e1ba45 100644
--- a/libgomp/libgomp-plugin.h
+++ b/libgomp/libgomp-plugin.h
@@ -112,7 +112,7 @@ extern void GOMP_OFFLOAD_openacc_exec (void (*) (void *), size_t, void **,
        void **, unsigned *, void *);
 extern void *GOMP_OFFLOAD_openacc_create_thread_data (int);
 extern void GOMP_OFFLOAD_openacc_destroy_thread_data (void *);
-extern struct goacc_asyncqueue *GOMP_OFFLOAD_openacc_async_construct (void);
+extern struct goacc_asyncqueue *GOMP_OFFLOAD_openacc_async_construct (int);
 extern bool GOMP_OFFLOAD_openacc_async_destruct (struct goacc_asyncqueue *);
 extern int GOMP_OFFLOAD_openacc_async_test (struct goacc_asyncqueue *);
 extern bool GOMP_OFFLOAD_openacc_async_synchronize (struct goacc_asyncqueue *);
diff --git a/libgomp/oacc-async.c b/libgomp/oacc-async.c
index 1760e8c90c6..2b24ae7adc2 100644
--- a/libgomp/oacc-async.c
+++ b/libgomp/oacc-async.c
@@ -100,7 +100,8 @@ lookup_goacc_asyncqueue (struct goacc_thread *thr, bool create, int async)
 
   if (!dev->openacc.async.asyncqueue[async])
     {
-      dev->openacc.async.asyncqueue[async] = dev->openacc.async.construct_func ();
+      dev->openacc.async.asyncqueue[async]
+ = dev->openacc.async.construct_func (dev->target_id);
 
       if (!dev->openacc.async.asyncqueue[async])
  {
diff --git a/libgomp/oacc-host.c b/libgomp/oacc-host.c
index 12299aee65d..cbcac9bf7b3 100644
--- a/libgomp/oacc-host.c
+++ b/libgomp/oacc-host.c
@@ -211,7 +211,7 @@ host_openacc_async_queue_callback (struct goacc_asyncqueue *aq
 }
 
 static struct goacc_asyncqueue *
-host_openacc_async_construct (void)
+host_openacc_async_construct (int device __attribute__((unused)))
 {
   /* Non-NULL 0xffff... value as opaque dummy.  */
   return (struct goacc_asyncqueue *) -1;
diff --git a/libgomp/plugin/plugin-nvptx.c b/libgomp/plugin/plugin-nvptx.c
index 9e088612b44..911d0f66a6e 100644
--- a/libgomp/plugin/plugin-nvptx.c
+++ b/libgomp/plugin/plugin-nvptx.c
@@ -1575,7 +1575,7 @@ GOMP_OFFLOAD_openacc_cuda_set_stream (struct goacc_asyncqueue *aq, void *stream)
 }
 
 struct goacc_asyncqueue *
-GOMP_OFFLOAD_openacc_async_construct (void)
+GOMP_OFFLOAD_openacc_async_construct (int device __attribute__((unused)))
 {
   CUstream stream = NULL;
   CUDA_CALL_ERET (NULL, cuStreamCreate, &stream, CU_STREAM_DEFAULT);
Reply | Threaded
Open this post in threaded view
|

[PATCH 4/7 libgomp,amdgcn] GCN libgomp port

Andrew Stubbs-4
In reply to this post by Andrew Stubbs-4
This patch contributes a libgomp implementation for AMD GCN, minus the
plugin which is later in this series.

GCN has been allocated ID number "8", even though devices "6" and "7"
are no longer present in every place where the IDs exist (they were HSA
and Intel MIC).

Most of these changes are simply based on the model already defined for
NVPTX, adapted for GCN as appropriate.  It is assumed that the "accel"
patch has already been applied, with the files that did not need to be
adjusted.

The "oacc-target.c" file is new.  This file allows new target-specific
symbols to be added to libgomp.  I couldn't find an existing way to do
this without adding a new top-level file also, to there's an empty
placeholder also.  (The OG9 branch has this symbol in libgcc, but that
seems wrong.)

OK to commit?

Thanks

Andrew


2019-11-12  Andrew Stubbs  <[hidden email]>

        include/
        * gomp-constants.h (GOMP_DEVICE_GCN): Define.
        (GOMP_VERSION_GCN): Define.

        libgomp/
        * Makefile.am (libgomp_la_SOURCES): Add oacc-target.c.
        * Makefile.in: Regenerate.
        * config.h.in (PLUGIN_GCN): Add new undef.
        * config/accel/openacc.f90 (acc_device_gcn): New parameter.
        * config/gcn/affinity-fmt.c: New file.
        * config/gcn/bar.c: New file.
        * config/gcn/bar.h: New file.
        * config/gcn/doacross.h: New file.
        * config/gcn/icv-device.c: New file.
        * config/gcn/oacc-target.c: New file.
        * config/gcn/simple-bar.h: New file.
        * config/gcn/target.c: New file.
        * config/gcn/task.c: New file.
        * config/gcn/team.c: New file.
        * config/gcn/time.c: New file.
        * configure.ac: Add amdgcn*-*-*.
        * configure: Regenerate.
        * configure.tgt: Add amdgcn*-*-*.
        * libgomp-plugin.h (offload_target_type): Add OFFLOAD_TARGET_TYPE_GCN.
        * libgomp.h (gcn_thrs): Add amdgcn variant.
        (set_gcn_thrs): Likewise.
        (gomp_thread): Likewise.
        * oacc-int.h (goacc_thread): Likewise.
        * oacc-target.c: New file.
        * openacc.f90 (acc_device_gcn): New parameter.
        * openacc.h (acc_device_t): Add acc_device_gcn.
        * team.c (gomp_free_pool_helper): Add amdgcn support.
---
 include/gomp-constants.h          |   2 +
 libgomp/Makefile.am               |   2 +-
 libgomp/config.h.in               |   3 +
 libgomp/config/accel/openacc.f90  |   1 +
 libgomp/config/gcn/affinity-fmt.c |  51 +++++++
 libgomp/config/gcn/bar.c          | 232 ++++++++++++++++++++++++++++++
 libgomp/config/gcn/bar.h          | 168 ++++++++++++++++++++++
 libgomp/config/gcn/doacross.h     |  58 ++++++++
 libgomp/config/gcn/icv-device.c   |  72 ++++++++++
 libgomp/config/gcn/oacc-target.c  |  31 ++++
 libgomp/config/gcn/simple-bar.h   |  61 ++++++++
 libgomp/config/gcn/target.c       |  67 +++++++++
 libgomp/config/gcn/task.c         |  39 +++++
 libgomp/config/gcn/team.c         | 202 ++++++++++++++++++++++++++
 libgomp/config/gcn/time.c         |  52 +++++++
 libgomp/configure.ac              |   2 +-
 libgomp/configure.tgt             |   4 +
 libgomp/libgomp-plugin.h          |   3 +-
 libgomp/libgomp.h                 |  18 +++
 libgomp/oacc-int.h                |   9 +-
 libgomp/oacc-target.c             |   1 +
 libgomp/openacc.f90               |   1 +
 libgomp/openacc.h                 |   1 +
 libgomp/team.c                    |   3 +
 26 files changed, 1189 insertions(+), 16 deletions(-)
 create mode 100644 libgomp/config/gcn/affinity-fmt.c
 create mode 100644 libgomp/config/gcn/bar.c
 create mode 100644 libgomp/config/gcn/bar.h
 create mode 100644 libgomp/config/gcn/doacross.h
 create mode 100644 libgomp/config/gcn/icv-device.c
 create mode 100644 libgomp/config/gcn/oacc-target.c
 create mode 100644 libgomp/config/gcn/simple-bar.h
 create mode 100644 libgomp/config/gcn/target.c
 create mode 100644 libgomp/config/gcn/task.c
 create mode 100644 libgomp/config/gcn/team.c
 create mode 100644 libgomp/config/gcn/time.c
 create mode 100644 libgomp/oacc-target.c


diff --git a/include/gomp-constants.h b/include/gomp-constants.h
index 82e9094c934..9e356cdfeec 100644
--- a/include/gomp-constants.h
+++ b/include/gomp-constants.h
@@ -174,6 +174,7 @@ enum gomp_map_kind
 #define GOMP_DEVICE_NVIDIA_PTX 5
 #define GOMP_DEVICE_INTEL_MIC 6
 #define GOMP_DEVICE_HSA 7
+#define GOMP_DEVICE_GCN 8
 
 #define GOMP_DEVICE_ICV -1
 #define GOMP_DEVICE_HOST_FALLBACK -2
@@ -215,6 +216,7 @@ enum gomp_map_kind
 #define GOMP_VERSION_NVIDIA_PTX 1
 #define GOMP_VERSION_INTEL_MIC 0
 #define GOMP_VERSION_HSA 0
+#define GOMP_VERSION_GCN 1
 
 #define GOMP_VERSION_PACK(LIB, DEV) (((LIB) << 16) | (DEV))
 #define GOMP_VERSION_LIB(PACK) (((PACK) >> 16) & 0xffff)
diff --git a/libgomp/Makefile.am b/libgomp/Makefile.am
index 7d36343a4be..669b9e4defd 100644
--- a/libgomp/Makefile.am
+++ b/libgomp/Makefile.am
@@ -65,7 +65,7 @@ libgomp_la_SOURCES = alloc.c atomic.c barrier.c critical.c env.c error.c \
  proc.c sem.c bar.c ptrlock.c time.c fortran.c affinity.c target.c \
  splay-tree.c libgomp-plugin.c oacc-parallel.c oacc-host.c oacc-init.c \
  oacc-mem.c oacc-async.c oacc-plugin.c oacc-cuda.c priority_queue.c \
- affinity-fmt.c teams.c oacc-profiling.c
+ affinity-fmt.c teams.c oacc-profiling.c oacc-target.c
 
 include $(top_srcdir)/plugin/Makefrag.am
 
diff --git a/libgomp/config.h.in b/libgomp/config.h.in
index ceb062fcb4c..2d50fcd5c1a 100644
--- a/libgomp/config.h.in
+++ b/libgomp/config.h.in
@@ -170,6 +170,9 @@
 /* Define to the version of this package. */
 #undef PACKAGE_VERSION
 
+/* Define to 1 if the GCN plugin is built, 0 if not. */
+#undef PLUGIN_GCN
+
 /* Define to 1 if the HSA plugin is built, 0 if not. */
 #undef PLUGIN_HSA
 
diff --git a/libgomp/config/accel/openacc.f90 b/libgomp/config/accel/openacc.f90
index a7f690e1572..6a8c5e9cb3d 100644
--- a/libgomp/config/accel/openacc.f90
+++ b/libgomp/config/accel/openacc.f90
@@ -51,6 +51,7 @@ module openacc_kinds
   ! integer (acc_device_kind), parameter :: acc_device_host_nonshm = 3 removed.
   integer (acc_device_kind), parameter :: acc_device_not_host = 4
   integer (acc_device_kind), parameter :: acc_device_nvidia = 5
+  integer (acc_device_kind), parameter :: acc_device_gcn = 8
 
 end module
 
diff --git a/libgomp/config/gcn/affinity-fmt.c b/libgomp/config/gcn/affinity-fmt.c
new file mode 100644
index 00000000000..3585f414460
--- /dev/null
+++ b/libgomp/config/gcn/affinity-fmt.c
@@ -0,0 +1,51 @@
+/* Copyright (C) 2018-2019 Free Software Foundation, Inc.
+
+   This file is part of the GNU Offloading and Multi Processing Library
+   (libgomp).
+
+   Libgomp is free software; you can redistribute it and/or modify it
+   under the terms of the GNU General Public License as published by
+   the Free Software Foundation; either version 3, or (at your option)
+   any later version.
+
+   Libgomp is distributed in the hope that it will be useful, but WITHOUT ANY
+   WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS
+   FOR A PARTICULAR PURPOSE.  See the GNU General Public License for
+   more details.
+
+   Under Section 7 of GPL version 3, you are granted additional
+   permissions described in the GCC Runtime Library Exception, version
+   3.1, as published by the Free Software Foundation.
+
+   You should have received a copy of the GNU General Public License and
+   a copy of the GCC Runtime Library Exception along with this program;
+   see the files COPYING3 and COPYING.RUNTIME respectively.  If not, see
+   <http://www.gnu.org/licenses/>.  */
+
+#include "libgomp.h"
+#include <string.h>
+#include <stdio.h>
+#include <stdlib.h>
+#ifdef HAVE_UNISTD_H
+#include <unistd.h>
+#endif
+#ifdef HAVE_INTTYPES_H
+# include <inttypes.h>  /* For PRIx64.  */
+#endif
+#ifdef HAVE_UNAME
+#include <sys/utsname.h>
+#endif
+
+/* The HAVE_GETPID and HAVE_GETHOSTNAME configure tests are passing for nvptx,
+   while the nvptx newlib implementation does not support those functions.
+   Override the configure test results here.  */
+#undef HAVE_GETPID
+#undef HAVE_GETHOSTNAME
+
+/* The GCN newlib implementation does not support fwrite, but it does support
+   write.  Map fwrite to write.  */
+#undef fwrite
+#define fwrite(ptr, size, nmemb, stream) write (1, (ptr), (nmemb) * (size))
+
+#include "../../affinity-fmt.c"
+
diff --git a/libgomp/config/gcn/bar.c b/libgomp/config/gcn/bar.c
new file mode 100644
index 00000000000..fb709be26ce
--- /dev/null
+++ b/libgomp/config/gcn/bar.c
@@ -0,0 +1,232 @@
+/* Copyright (C) 2015-2019 Free Software Foundation, Inc.
+   Contributed by Mentor Embedded.
+
+   This file is part of the GNU Offloading and Multi Processing Library
+   (libgomp).
+
+   Libgomp is free software; you can redistribute it and/or modify it
+   under the terms of the GNU General Public License as published by
+   the Free Software Foundation; either version 3, or (at your option)
+   any later version.
+
+   Libgomp is distributed in the hope that it will be useful, but WITHOUT ANY
+   WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS
+   FOR A PARTICULAR PURPOSE.  See the GNU General Public License for
+   more details.
+
+   Under Section 7 of GPL version 3, you are granted additional
+   permissions described in the GCC Runtime Library Exception, version
+   3.1, as published by the Free Software Foundation.
+
+   You should have received a copy of the GNU General Public License and
+   a copy of the GCC Runtime Library Exception along with this program;
+   see the files COPYING3 and COPYING.RUNTIME respectively.  If not, see
+   <http://www.gnu.org/licenses/>.  */
+
+/* This is an AMD GCN specific implementation of a barrier synchronization
+   mechanism for libgomp.  This type is private to the library.  This
+   implementation uses atomic instructions and s_barrier instruction.  It
+   uses MEMMODEL_RELAXED here because barriers are within workgroups and
+   therefore don't need to flush caches.  */
+
+#include <limits.h>
+#include "libgomp.h"
+
+
+void
+gomp_barrier_wait_end (gomp_barrier_t *bar, gomp_barrier_state_t state)
+{
+  if (__builtin_expect (state & BAR_WAS_LAST, 0))
+    {
+      /* Next time we'll be awaiting TOTAL threads again.  */
+      bar->awaited = bar->total;
+      __atomic_store_n (&bar->generation, bar->generation + BAR_INCR,
+ MEMMODEL_RELAXED);
+    }
+  asm ("s_barrier" ::: "memory");
+}
+
+void
+gomp_barrier_wait (gomp_barrier_t *bar)
+{
+  gomp_barrier_wait_end (bar, gomp_barrier_wait_start (bar));
+}
+
+/* Like gomp_barrier_wait, except that if the encountering thread
+   is not the last one to hit the barrier, it returns immediately.
+   The intended usage is that a thread which intends to gomp_barrier_destroy
+   this barrier calls gomp_barrier_wait, while all other threads
+   call gomp_barrier_wait_last.  When gomp_barrier_wait returns,
+   the barrier can be safely destroyed.  */
+
+void
+gomp_barrier_wait_last (gomp_barrier_t *bar)
+{
+  /* Deferring to gomp_barrier_wait does not use the optimization opportunity
+     allowed by the interface contract for all-but-last participants.  The
+     original implementation in config/linux/bar.c handles this better.  */
+  gomp_barrier_wait (bar);
+}
+
+void
+gomp_team_barrier_wake (gomp_barrier_t *bar, int count)
+{
+  asm ("s_barrier" ::: "memory");
+}
+
+void
+gomp_team_barrier_wait_end (gomp_barrier_t *bar, gomp_barrier_state_t state)
+{
+  unsigned int generation, gen;
+
+  if (__builtin_expect (state & BAR_WAS_LAST, 0))
+    {
+      /* Next time we'll be awaiting TOTAL threads again.  */
+      struct gomp_thread *thr = gomp_thread ();
+      struct gomp_team *team = thr->ts.team;
+
+      bar->awaited = bar->total;
+      team->work_share_cancelled = 0;
+      if (__builtin_expect (team->task_count, 0))
+ {
+  gomp_barrier_handle_tasks (state);
+  state &= ~BAR_WAS_LAST;
+ }
+      else
+ {
+  state &= ~BAR_CANCELLED;
+  state += BAR_INCR - BAR_WAS_LAST;
+  __atomic_store_n (&bar->generation, state, MEMMODEL_RELAXED);
+  asm ("s_barrier" ::: "memory");
+  return;
+ }
+    }
+
+  generation = state;
+  state &= ~BAR_CANCELLED;
+  int retry = 100;
+  do
+    {
+      if (retry-- == 0)
+ {
+  /* It really shouldn't happen that barriers get out of sync, but
+     if they do then this will loop until they realign, so we need
+     to avoid an infinite loop where the thread just isn't there.  */
+  const char msg[] = ("Barrier sync failed (another thread died?);"
+      " aborting.");
+  write (2, msg, sizeof (msg)-1);
+  abort();
+ }
+
+      asm ("s_barrier" ::: "memory");
+      gen = __atomic_load_n (&bar->generation, MEMMODEL_ACQUIRE);
+      if (__builtin_expect (gen & BAR_TASK_PENDING, 0))
+ {
+  gomp_barrier_handle_tasks (state);
+  gen = __atomic_load_n (&bar->generation, MEMMODEL_ACQUIRE);
+ }
+      generation |= gen & BAR_WAITING_FOR_TASK;
+    }
+  while (gen != state + BAR_INCR);
+}
+
+void
+gomp_team_barrier_wait (gomp_barrier_t *bar)
+{
+  gomp_team_barrier_wait_end (bar, gomp_barrier_wait_start (bar));
+}
+
+void
+gomp_team_barrier_wait_final (gomp_barrier_t *bar)
+{
+  gomp_barrier_state_t state = gomp_barrier_wait_final_start (bar);
+  if (__builtin_expect (state & BAR_WAS_LAST, 0))
+    bar->awaited_final = bar->total;
+  gomp_team_barrier_wait_end (bar, state);
+}
+
+bool
+gomp_team_barrier_wait_cancel_end (gomp_barrier_t *bar,
+   gomp_barrier_state_t state)
+{
+  unsigned int generation, gen;
+
+  if (__builtin_expect (state & BAR_WAS_LAST, 0))
+    {
+      /* Next time we'll be awaiting TOTAL threads again.  */
+      /* BAR_CANCELLED should never be set in state here, because
+ cancellation means that at least one of the threads has been
+ cancelled, thus on a cancellable barrier we should never see
+ all threads to arrive.  */
+      struct gomp_thread *thr = gomp_thread ();
+      struct gomp_team *team = thr->ts.team;
+
+      bar->awaited = bar->total;
+      team->work_share_cancelled = 0;
+      if (__builtin_expect (team->task_count, 0))
+ {
+  gomp_barrier_handle_tasks (state);
+  state &= ~BAR_WAS_LAST;
+ }
+      else
+ {
+  state += BAR_INCR - BAR_WAS_LAST;
+  __atomic_store_n (&bar->generation, state, MEMMODEL_RELAXED);
+  asm ("s_barrier" ::: "memory");
+  return false;
+ }
+    }
+
+  if (__builtin_expect (state & BAR_CANCELLED, 0))
+    return true;
+
+  generation = state;
+  int retry = 100;
+  do
+    {
+      if (retry-- == 0)
+ {
+  /* It really shouldn't happen that barriers get out of sync, but
+     if they do then this will loop until they realign, so we need
+     to avoid an infinite loop where the thread just isn't there.  */
+  const char msg[] = ("Barrier sync failed (another thread died?);"
+      " aborting.");
+  write (2, msg, sizeof (msg)-1);
+  abort();
+ }
+
+      asm ("s_barrier" ::: "memory");
+      gen = __atomic_load_n (&bar->generation, MEMMODEL_RELAXED);
+      if (__builtin_expect (gen & BAR_CANCELLED, 0))
+ return true;
+      if (__builtin_expect (gen & BAR_TASK_PENDING, 0))
+ {
+  gomp_barrier_handle_tasks (state);
+  gen = __atomic_load_n (&bar->generation, MEMMODEL_RELAXED);
+ }
+      generation |= gen & BAR_WAITING_FOR_TASK;
+    }
+  while (gen != state + BAR_INCR);
+
+  return false;
+}
+
+bool
+gomp_team_barrier_wait_cancel (gomp_barrier_t *bar)
+{
+  return gomp_team_barrier_wait_cancel_end (bar, gomp_barrier_wait_start (bar));
+}
+
+void
+gomp_team_barrier_cancel (struct gomp_team *team)
+{
+  gomp_mutex_lock (&team->task_lock);
+  if (team->barrier.generation & BAR_CANCELLED)
+    {
+      gomp_mutex_unlock (&team->task_lock);
+      return;
+    }
+  team->barrier.generation |= BAR_CANCELLED;
+  gomp_mutex_unlock (&team->task_lock);
+  gomp_team_barrier_wake (&team->barrier, INT_MAX);
+}
diff --git a/libgomp/config/gcn/bar.h b/libgomp/config/gcn/bar.h
new file mode 100644
index 00000000000..ec8851ba078
--- /dev/null
+++ b/libgomp/config/gcn/bar.h
@@ -0,0 +1,168 @@
+/* Copyright (C) 2015-2019 Free Software Foundation, Inc.
+   Contributed by Mentor Embedded.
+
+   This file is part of the GNU Offloading and Multi Processing Library
+   (libgomp).
+
+   Libgomp is free software; you can redistribute it and/or modify it
+   under the terms of the GNU General Public License as published by
+   the Free Software Foundation; either version 3, or (at your option)
+   any later version.
+
+   Libgomp is distributed in the hope that it will be useful, but WITHOUT ANY
+   WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS
+   FOR A PARTICULAR PURPOSE.  See the GNU General Public License for
+   more details.
+
+   Under Section 7 of GPL version 3, you are granted additional
+   permissions described in the GCC Runtime Library Exception, version
+   3.1, as published by the Free Software Foundation.
+
+   You should have received a copy of the GNU General Public License and
+   a copy of the GCC Runtime Library Exception along with this program;
+   see the files COPYING3 and COPYING.RUNTIME respectively.  If not, see
+   <http://www.gnu.org/licenses/>.  */
+
+/* This is an AMD GCN specific implementation of a barrier synchronization
+   mechanism for libgomp.  This type is private to the library.  This
+   implementation uses atomic instructions and s_barrier instruction.  It
+   uses MEMMODEL_RELAXED here because barriers are within workgroups and
+   therefore don't need to flush caches.  */
+
+#ifndef GOMP_BARRIER_H
+#define GOMP_BARRIER_H 1
+
+#include "mutex.h"
+
+typedef struct
+{
+  unsigned total;
+  unsigned generation;
+  unsigned awaited;
+  unsigned awaited_final;
+} gomp_barrier_t;
+
+typedef unsigned int gomp_barrier_state_t;
+
+/* The generation field contains a counter in the high bits, with a few
+   low bits dedicated to flags.  Note that TASK_PENDING and WAS_LAST can
+   share space because WAS_LAST is never stored back to generation.  */
+#define BAR_TASK_PENDING 1
+#define BAR_WAS_LAST 1
+#define BAR_WAITING_FOR_TASK 2
+#define BAR_CANCELLED 4
+#define BAR_INCR 8
+
+static inline void gomp_barrier_init (gomp_barrier_t *bar, unsigned count)
+{
+  bar->total = count;
+  bar->awaited = count;
+  bar->awaited_final = count;
+  bar->generation = 0;
+}
+
+static inline void gomp_barrier_reinit (gomp_barrier_t *bar, unsigned count)
+{
+  __atomic_add_fetch (&bar->awaited, count - bar->total, MEMMODEL_RELAXED);
+  bar->total = count;
+}
+
+static inline void gomp_barrier_destroy (gomp_barrier_t *bar)
+{
+}
+
+extern void gomp_barrier_wait (gomp_barrier_t *);
+extern void gomp_barrier_wait_last (gomp_barrier_t *);
+extern void gomp_barrier_wait_end (gomp_barrier_t *, gomp_barrier_state_t);
+extern void gomp_team_barrier_wait (gomp_barrier_t *);
+extern void gomp_team_barrier_wait_final (gomp_barrier_t *);
+extern void gomp_team_barrier_wait_end (gomp_barrier_t *,
+ gomp_barrier_state_t);
+extern bool gomp_team_barrier_wait_cancel (gomp_barrier_t *);
+extern bool gomp_team_barrier_wait_cancel_end (gomp_barrier_t *,
+       gomp_barrier_state_t);
+extern void gomp_team_barrier_wake (gomp_barrier_t *, int);
+struct gomp_team;
+extern void gomp_team_barrier_cancel (struct gomp_team *);
+
+static inline gomp_barrier_state_t
+gomp_barrier_wait_start (gomp_barrier_t *bar)
+{
+  unsigned int ret = __atomic_load_n (&bar->generation, MEMMODEL_RELAXED);
+  ret &= -BAR_INCR | BAR_CANCELLED;
+  /* A memory barrier is needed before exiting from the various forms
+     of gomp_barrier_wait, to satisfy OpenMP API version 3.1 section
+     2.8.6 flush Construct, which says there is an implicit flush during
+     a barrier region.  This is a convenient place to add the barrier,
+     so we use MEMMODEL_ACQ_REL here rather than MEMMODEL_ACQUIRE.  */
+  if (__atomic_add_fetch (&bar->awaited, -1, MEMMODEL_RELAXED) == 0)
+    ret |= BAR_WAS_LAST;
+  return ret;
+}
+
+static inline gomp_barrier_state_t
+gomp_barrier_wait_cancel_start (gomp_barrier_t *bar)
+{
+  return gomp_barrier_wait_start (bar);
+}
+
+/* This is like gomp_barrier_wait_start, except it decrements
+   bar->awaited_final rather than bar->awaited and should be used
+   for the gomp_team_end barrier only.  */
+static inline gomp_barrier_state_t
+gomp_barrier_wait_final_start (gomp_barrier_t *bar)
+{
+  unsigned int ret = __atomic_load_n (&bar->generation, MEMMODEL_RELAXED);
+  ret &= -BAR_INCR | BAR_CANCELLED;
+  /* See above gomp_barrier_wait_start comment.  */
+  if (__atomic_add_fetch (&bar->awaited_final, -1, MEMMODEL_RELAXED) == 0)
+    ret |= BAR_WAS_LAST;
+  return ret;
+}
+
+static inline bool
+gomp_barrier_last_thread (gomp_barrier_state_t state)
+{
+  return state & BAR_WAS_LAST;
+}
+
+/* All the inlines below must be called with team->task_lock
+   held.  */
+
+static inline void
+gomp_team_barrier_set_task_pending (gomp_barrier_t *bar)
+{
+  bar->generation |= BAR_TASK_PENDING;
+}
+
+static inline void
+gomp_team_barrier_clear_task_pending (gomp_barrier_t *bar)
+{
+  bar->generation &= ~BAR_TASK_PENDING;
+}
+
+static inline void
+gomp_team_barrier_set_waiting_for_tasks (gomp_barrier_t *bar)
+{
+  bar->generation |= BAR_WAITING_FOR_TASK;
+}
+
+static inline bool
+gomp_team_barrier_waiting_for_tasks (gomp_barrier_t *bar)
+{
+  return (bar->generation & BAR_WAITING_FOR_TASK) != 0;
+}
+
+static inline bool
+gomp_team_barrier_cancelled (gomp_barrier_t *bar)
+{
+  return __builtin_expect ((bar->generation & BAR_CANCELLED) != 0, 0);
+}
+
+static inline void
+gomp_team_barrier_done (gomp_barrier_t *bar, gomp_barrier_state_t state)
+{
+  bar->generation = (state & -BAR_INCR) + BAR_INCR;
+}
+
+#endif /* GOMP_BARRIER_H */
diff --git a/libgomp/config/gcn/doacross.h b/libgomp/config/gcn/doacross.h
new file mode 100644
index 00000000000..2bff18ae1a8
--- /dev/null
+++ b/libgomp/config/gcn/doacross.h
@@ -0,0 +1,58 @@
+/* Copyright (C) 2015-2019 Free Software Foundation, Inc.
+   Contributed by Mentor Embedded.
+
+   This file is part of the GNU Offloading and Multi Processing Library
+   (libgomp).
+
+   Libgomp is free software; you can redistribute it and/or modify it
+   under the terms of the GNU General Public License as published by
+   the Free Software Foundation; either version 3, or (at your option)
+   any later version.
+
+   Libgomp is distributed in the hope that it will be useful, but WITHOUT ANY
+   WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS
+   FOR A PARTICULAR PURPOSE.  See the GNU General Public License for
+   more details.
+
+   Under Section 7 of GPL version 3, you are granted additional
+   permissions described in the GCC Runtime Library Exception, version
+   3.1, as published by the Free Software Foundation.
+
+   You should have received a copy of the GNU General Public License and
+   a copy of the GCC Runtime Library Exception along with this program;
+   see the files COPYING3 and COPYING.RUNTIME respectively.  If not, see
+   <http://www.gnu.org/licenses/>.  */
+
+/* This is the AMD GCN implementation of doacross spinning.  */
+
+#ifndef GOMP_DOACROSS_H
+#define GOMP_DOACROSS_H 1
+
+#include "libgomp.h"
+
+static inline int
+cpu_relax (void)
+{
+  /* This can be implemented as just a memory barrier, but a sleep seems
+     like it should allow the wavefront to yield (maybe?)
+     Use the shortest possible sleep time of 1*64 cycles.  */
+  asm volatile ("s_sleep\t1" ::: "memory");
+  return 0;
+}
+
+static inline void doacross_spin (unsigned long *addr, unsigned long expected,
+  unsigned long cur)
+{
+  /* Prevent compiler from optimizing based on bounds of containing object.  */
+  asm ("" : "+r" (addr));
+  do
+    {
+       /* An alternative implementation might use s_setprio to lower the
+  priority temporarily, and then restore it after.  */
+      int i = cpu_relax ();
+      cur = addr[i];
+    }
+  while (cur <= expected);
+}
+
+#endif /* GOMP_DOACROSS_H */
diff --git a/libgomp/config/gcn/icv-device.c b/libgomp/config/gcn/icv-device.c
new file mode 100644
index 00000000000..cbb9dfa1133
--- /dev/null
+++ b/libgomp/config/gcn/icv-device.c
@@ -0,0 +1,72 @@
+/* Copyright (C) 2015-2019 Free Software Foundation, Inc.
+   Contributed by Mentor Embedded.
+
+   This file is part of the GNU Offloading and Multi Processing Library
+   (libgomp).
+
+   Libgomp is free software; you can redistribute it and/or modify it
+   under the terms of the GNU General Public License as published by
+   the Free Software Foundation; either version 3, or (at your option)
+   any later version.
+
+   Libgomp is distributed in the hope that it will be useful, but WITHOUT ANY
+   WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS
+   FOR A PARTICULAR PURPOSE.  See the GNU General Public License for
+   more details.
+
+   Under Section 7 of GPL version 3, you are granted additional
+   permissions described in the GCC Runtime Library Exception, version
+   3.1, as published by the Free Software Foundation.
+
+   You should have received a copy of the GNU General Public License and
+   a copy of the GCC Runtime Library Exception along with this program;
+   see the files COPYING3 and COPYING.RUNTIME respectively.  If not, see
+   <http://www.gnu.org/licenses/>.  */
+
+/* This file defines OpenMP API entry points that accelerator targets are
+   expected to replace.  */
+
+#include "libgomp.h"
+
+void
+omp_set_default_device (int device_num __attribute__((unused)))
+{
+}
+
+int
+omp_get_default_device (void)
+{
+  return 0;
+}
+
+int
+omp_get_num_devices (void)
+{
+  return 0;
+}
+
+int
+omp_get_num_teams (void)
+{
+  return gomp_num_teams_var + 1;
+}
+
+int __attribute__ ((__optimize__ ("O2")))
+omp_get_team_num (void)
+{
+  return __builtin_gcn_dim_pos (0);
+}
+
+int
+omp_is_initial_device (void)
+{
+  /* AMD GCN is an accelerator-only target.  */
+  return 0;
+}
+
+ialias (omp_set_default_device)
+ialias (omp_get_default_device)
+ialias (omp_get_num_devices)
+ialias (omp_get_num_teams)
+ialias (omp_get_team_num)
+ialias (omp_is_initial_device)
diff --git a/libgomp/config/gcn/oacc-target.c b/libgomp/config/gcn/oacc-target.c
new file mode 100644
index 00000000000..bdcc9153d96
--- /dev/null
+++ b/libgomp/config/gcn/oacc-target.c
@@ -0,0 +1,31 @@
+/* Oversized reductions lock variable
+   Copyright (C) 2017-2019 Free Software Foundation, Inc.
+   Contributed by Mentor Graphics.
+
+   This file is part of the GNU Offloading and Multi Processing Library
+   (libgomp).
+
+   Libgomp is free software; you can redistribute it and/or modify it
+   under the terms of the GNU General Public License as published by
+   the Free Software Foundation; either version 3, or (at your option)
+   any later version.
+
+   Libgomp is distributed in the hope that it will be useful, but WITHOUT ANY
+   WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS
+   FOR A PARTICULAR PURPOSE.  See the GNU General Public License for
+   more details.
+
+   Under Section 7 of GPL version 3, you are granted additional
+   permissions described in the GCC Runtime Library Exception, version
+   3.1, as published by the Free Software Foundation.
+
+   You should have received a copy of the GNU General Public License and
+   a copy of the GCC Runtime Library Exception along with this program;
+   see the files COPYING3 and COPYING.RUNTIME respectively.  If not, see
+   <http://www.gnu.org/licenses/>.  */
+
+/* We use a global lock variable for reductions on objects larger than
+   64 bits.  Until and unless proven that lock contention for
+   different reductions is a problem, a single lock will suffice.  */
+
+unsigned volatile __reduction_lock = 0;
diff --git a/libgomp/config/gcn/simple-bar.h b/libgomp/config/gcn/simple-bar.h
new file mode 100644
index 00000000000..802e0f5c301
--- /dev/null
+++ b/libgomp/config/gcn/simple-bar.h
@@ -0,0 +1,61 @@
+/* Copyright (C) 2015-2019 Free Software Foundation, Inc.
+   Contributed by Mentor Embedded.
+
+   This file is part of the GNU Offloading and Multi Processing Library
+   (libgomp).
+
+   Libgomp is free software; you can redistribute it and/or modify it
+   under the terms of the GNU General Public License as published by
+   the Free Software Foundation; either version 3, or (at your option)
+   any later version.
+
+   Libgomp is distributed in the hope that it will be useful, but WITHOUT ANY
+   WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS
+   FOR A PARTICULAR PURPOSE.  See the GNU General Public License for
+   more details.
+
+   Under Section 7 of GPL version 3, you are granted additional
+   permissions described in the GCC Runtime Library Exception, version
+   3.1, as published by the Free Software Foundation.
+
+   You should have received a copy of the GNU General Public License and
+   a copy of the GCC Runtime Library Exception along with this program;
+   see the files COPYING3 and COPYING.RUNTIME respectively.  If not, see
+   <http://www.gnu.org/licenses/>.  */
+
+/* This is a simplified barrier that is suitable for thread pool
+   synchronizaton.  Only a subset of full barrier API (bar.h) is exposed.
+   Here in the AMD GCN-specific implementation, we expect that thread pool
+   corresponds to the wavefronts within a work group.  */
+
+#ifndef GOMP_SIMPLE_BARRIER_H
+#define GOMP_SIMPLE_BARRIER_H 1
+
+/* AMD GCN has no use for this type.  */
+typedef int gomp_simple_barrier_t;
+
+/* GCN barriers block all wavefronts, so the count is not interesting.  */
+static inline void
+gomp_simple_barrier_init (gomp_simple_barrier_t *bar, unsigned count)
+{
+}
+
+static inline void
+gomp_simple_barrier_destroy (gomp_simple_barrier_t *bar)
+{
+}
+
+static inline void
+gomp_simple_barrier_wait (gomp_simple_barrier_t *bar)
+{
+  asm volatile ("s_barrier" ::: "memory");
+}
+
+static inline void
+gomp_simple_barrier_wait_last (gomp_simple_barrier_t *bar)
+{
+  /* GCN has no way to signal a barrier without waiting.  */
+  asm volatile ("s_barrier" ::: "memory");
+}
+
+#endif /* GOMP_SIMPLE_BARRIER_H */
diff --git a/libgomp/config/gcn/target.c b/libgomp/config/gcn/target.c
new file mode 100644
index 00000000000..db00551e695
--- /dev/null
+++ b/libgomp/config/gcn/target.c
@@ -0,0 +1,67 @@
+/* Copyright (C) 2017-2019 Free Software Foundation, Inc.
+   Contributed by Mentor Embedded.
+
+   This file is part of the GNU Offloading and Multi Processing Library
+   (libgomp).
+
+   Libgomp is free software; you can redistribute it and/or modify it
+   under the terms of the GNU General Public License as published by
+   the Free Software Foundation; either version 3, or (at your option)
+   any later version.
+
+   Libgomp is distributed in the hope that it will be useful, but WITHOUT ANY
+   WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS
+   FOR A PARTICULAR PURPOSE.  See the GNU General Public License for
+   more details.
+
+   Under Section 7 of GPL version 3, you are granted additional
+   permissions described in the GCC Runtime Library Exception, version
+   3.1, as published by the Free Software Foundation.
+
+   You should have received a copy of the GNU General Public License and
+   a copy of the GCC Runtime Library Exception along with this program;
+   see the files COPYING3 and COPYING.RUNTIME respectively.  If not, see
+   <http://www.gnu.org/licenses/>.  */
+
+#include "libgomp.h"
+#include <limits.h>
+
+void
+GOMP_teams (unsigned int num_teams, unsigned int thread_limit)
+{
+  if (thread_limit)
+    {
+      struct gomp_task_icv *icv = gomp_icv (true);
+      icv->thread_limit_var
+ = thread_limit > INT_MAX ? UINT_MAX : thread_limit;
+    }
+  unsigned int num_workgroups, workgroup_id;
+  num_workgroups = __builtin_gcn_dim_size (0);
+  workgroup_id = __builtin_gcn_dim_pos (0);
+  if (!num_teams || num_teams >= num_workgroups)
+    num_teams = num_workgroups;
+  else if (workgroup_id >= num_teams)
+    {
+      gomp_free_thread (gcn_thrs ());
+      exit (0);
+    }
+  gomp_num_teams_var = num_teams - 1;
+}
+
+int
+omp_pause_resource (omp_pause_resource_t kind, int device_num)
+{
+  (void) kind;
+  (void) device_num;
+  return -1;
+}
+
+int
+omp_pause_resource_all (omp_pause_resource_t kind)
+{
+  (void) kind;
+  return -1;
+}
+
+ialias (omp_pause_resource)
+ialias (omp_pause_resource_all)
diff --git a/libgomp/config/gcn/task.c b/libgomp/config/gcn/task.c
new file mode 100644
index 00000000000..a13565034b6
--- /dev/null
+++ b/libgomp/config/gcn/task.c
@@ -0,0 +1,39 @@
+/* Copyright (C) 2017-2019 Free Software Foundation, Inc.
+   Contributed by Mentor Embedded.
+
+   This file is part of the GNU Offloading and Multi Processing Library
+   (libgomp).
+
+   Libgomp is free software; you can redistribute it and/or modify it
+   under the terms of the GNU General Public License as published by
+   the Free Software Foundation; either version 3, or (at your option)
+   any later version.
+
+   Libgomp is distributed in the hope that it will be useful, but WITHOUT ANY
+   WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS
+   FOR A PARTICULAR PURPOSE.  See the GNU General Public License for
+   more details.
+
+   Under Section 7 of GPL version 3, you are granted additional
+   permissions described in the GCC Runtime Library Exception, version
+   3.1, as published by the Free Software Foundation.
+
+   You should have received a copy of the GNU General Public License and
+   a copy of the GCC Runtime Library Exception along with this program;
+   see the files COPYING3 and COPYING.RUNTIME respectively.  If not, see
+   <http://www.gnu.org/licenses/>.  */
+
+/* This file handles the maintainence of tasks in response to task
+   creation and termination.  */
+
+#include "libgomp.h"
+
+/* AMD GCN is an accelerator-only target, so this should never be called.  */
+
+bool
+gomp_target_task_fn (void *data)
+{
+  __builtin_unreachable ();
+}
+
+#include "../../task.c"
diff --git a/libgomp/config/gcn/team.c b/libgomp/config/gcn/team.c
new file mode 100644
index 00000000000..c566482bda2
--- /dev/null
+++ b/libgomp/config/gcn/team.c
@@ -0,0 +1,202 @@
+/* Copyright (C) 2017-2019 Free Software Foundation, Inc.
+   Contributed by Mentor Embedded.
+
+   This file is part of the GNU Offloading and Multi Processing Library
+   (libgomp).
+
+   Libgomp is free software; you can redistribute it and/or modify it
+   under the terms of the GNU General Public License as published by
+   the Free Software Foundation; either version 3, or (at your option)
+   any later version.
+
+   Libgomp is distributed in the hope that it will be useful, but WITHOUT ANY
+   WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS
+   FOR A PARTICULAR PURPOSE.  See the GNU General Public License for
+   more details.
+
+   Under Section 7 of GPL version 3, you are granted additional
+   permissions described in the GCC Runtime Library Exception, version
+   3.1, as published by the Free Software Foundation.
+
+   You should have received a copy of the GNU General Public License and
+   a copy of the GCC Runtime Library Exception along with this program;
+   see the files COPYING3 and COPYING.RUNTIME respectively.  If not, see
+   <http://www.gnu.org/licenses/>.  */
+
+/* This file handles maintainance of threads on AMD GCN.  */
+
+#include "libgomp.h"
+#include <stdlib.h>
+#include <string.h>
+
+static void gomp_thread_start (struct gomp_thread_pool *);
+
+/* This externally visible function handles target region entry.  It
+   sets up a per-team thread pool and transfers control by returning to
+   the kernel in the master thread or gomp_thread_start in other threads.
+
+   The name of this function is part of the interface with the compiler: for
+   each OpenMP kernel the compiler configures the stack, then calls here.
+
+   Likewise, gomp_gcn_exit_kernel is called during the kernel epilogue.  */
+
+void
+gomp_gcn_enter_kernel (void)
+{
+  int threadid = __builtin_gcn_dim_pos (1);
+
+  if (threadid == 0)
+    {
+      int numthreads = __builtin_gcn_dim_size (1);
+      int teamid = __builtin_gcn_dim_pos(0);
+
+      /* Set up the global state.
+ Every team will do this, but that should be harmless.  */
+      gomp_global_icv.nthreads_var = 16;
+      gomp_global_icv.thread_limit_var = numthreads;
+      /* Starting additional threads is not supported.  */
+      gomp_global_icv.dyn_var = true;
+
+      /* Allocate and initialize the team-local-storage data.  */
+      struct gomp_thread *thrs = gomp_malloc_cleared (sizeof (*thrs)
+      * numthreads);
+      set_gcn_thrs (thrs);
+
+      /* Allocate and initailize a pool of threads in the team.
+         The threads are already running, of course, we just need to manage
+         the communication between them.  */
+      struct gomp_thread_pool *pool = gomp_malloc (sizeof (*pool));
+      pool->threads = gomp_malloc (sizeof (void *) * numthreads);
+      for (int tid = 0; tid < numthreads; tid++)
+ pool->threads[tid] = &thrs[tid];
+      pool->threads_size = numthreads;
+      pool->threads_used = numthreads;
+      pool->threads_busy = 1;
+      pool->last_team = NULL;
+      gomp_simple_barrier_init (&pool->threads_dock, numthreads);
+      thrs->thread_pool = pool;
+
+      asm ("s_barrier" ::: "memory");
+      return;  /* Return to kernel.  */
+    }
+  else
+    {
+      asm ("s_barrier" ::: "memory");
+      gomp_thread_start (gcn_thrs ()[0].thread_pool);
+      /* gomp_thread_start does not return.  */
+    }
+}
+
+void
+gomp_gcn_exit_kernel (void)
+{
+  gomp_free_thread (gcn_thrs ());
+  free (gcn_thrs ());
+}
+
+/* This function contains the idle loop in which a thread waits
+   to be called up to become part of a team.  */
+
+static void
+gomp_thread_start (struct gomp_thread_pool *pool)
+{
+  struct gomp_thread *thr = gomp_thread ();
+
+  gomp_sem_init (&thr->release, 0);
+  thr->thread_pool = pool;
+
+  /* The loop exits only when "fn" is assigned "gomp_free_pool_helper",
+     which contains "s_endpgm", or an infinite no-op loop is
+     suspected (this happens when the thread master crashes).  */
+  int nul_limit = 99;
+  do
+    {
+      gomp_simple_barrier_wait (&pool->threads_dock);
+      if (!thr->fn)
+ {
+  if (nul_limit-- > 0)
+    continue;
+  else
+    {
+      const char msg[] = ("team master not responding;"
+  " slave thread aborting");
+      write (2, msg, sizeof (msg)-1);
+      abort();
+    }
+ }
+      thr->fn (thr->data);
+      thr->fn = NULL;
+
+      struct gomp_task *task = thr->task;
+      gomp_team_barrier_wait_final (&thr->ts.team->barrier);
+      gomp_finish_task (task);
+    }
+  while (1);
+}
+
+/* Launch a team.  */
+
+void
+gomp_team_start (void (*fn) (void *), void *data, unsigned nthreads,
+ unsigned flags, struct gomp_team *team,
+ struct gomp_taskgroup *taskgroup)
+{
+  struct gomp_thread *thr, *nthr;
+  struct gomp_task *task;
+  struct gomp_task_icv *icv;
+  struct gomp_thread_pool *pool;
+  unsigned long nthreads_var;
+
+  thr = gomp_thread ();
+  pool = thr->thread_pool;
+  task = thr->task;
+  icv = task ? &task->icv : &gomp_global_icv;
+
+  /* Always save the previous state, even if this isn't a nested team.
+     In particular, we should save any work share state from an outer
+     orphaned work share construct.  */
+  team->prev_ts = thr->ts;
+
+  thr->ts.team = team;
+  thr->ts.team_id = 0;
+  ++thr->ts.level;
+  if (nthreads > 1)
+    ++thr->ts.active_level;
+  thr->ts.work_share = &team->work_shares[0];
+  thr->ts.last_work_share = NULL;
+  thr->ts.single_count = 0;
+  thr->ts.static_trip = 0;
+  thr->task = &team->implicit_task[0];
+  nthreads_var = icv->nthreads_var;
+  gomp_init_task (thr->task, task, icv);
+  team->implicit_task[0].icv.nthreads_var = nthreads_var;
+  team->implicit_task[0].taskgroup = taskgroup;
+
+  if (nthreads == 1)
+    return;
+
+  /* Release existing idle threads.  */
+  for (unsigned i = 1; i < nthreads; ++i)
+    {
+      nthr = pool->threads[i];
+      nthr->ts.team = team;
+      nthr->ts.work_share = &team->work_shares[0];
+      nthr->ts.last_work_share = NULL;
+      nthr->ts.team_id = i;
+      nthr->ts.level = team->prev_ts.level + 1;
+      nthr->ts.active_level = thr->ts.active_level;
+      nthr->ts.single_count = 0;
+      nthr->ts.static_trip = 0;
+      nthr->task = &team->implicit_task[i];
+      gomp_init_task (nthr->task, task, icv);
+      team->implicit_task[i].icv.nthreads_var = nthreads_var;
+      team->implicit_task[i].taskgroup = taskgroup;
+      nthr->fn = fn;
+      nthr->data = data;
+      team->ordered_release[i] = &nthr->release;
+    }
+
+  gomp_simple_barrier_wait (&pool->threads_dock);
+}
+
+#include "../../team.c"
diff --git a/libgomp/config/gcn/time.c b/libgomp/config/gcn/time.c
new file mode 100644
index 00000000000..f189e55889c
--- /dev/null
+++ b/libgomp/config/gcn/time.c
@@ -0,0 +1,52 @@
+/* Copyright (C) 2015-2019 Free Software Foundation, Inc.
+   Contributed by Mentor Embedded.
+
+   This file is part of the GNU Offloading and Multi Processing Library
+   (libgomp).
+
+   Libgomp is free software; you can redistribute it and/or modify it
+   under the terms of the GNU General Public License as published by
+   the Free Software Foundation; either version 3, or (at your option)
+   any later version.
+
+   Libgomp is distributed in the hope that it will be useful, but WITHOUT ANY
+   WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS
+   FOR A PARTICULAR PURPOSE.  See the GNU General Public License for
+   more details.
+
+   Under Section 7 of GPL version 3, you are granted additional
+   permissions described in the GCC Runtime Library Exception, version
+   3.1, as published by the Free Software Foundation.
+
+   You should have received a copy of the GNU General Public License and
+   a copy of the GCC Runtime Library Exception along with this program;
+   see the files COPYING3 and COPYING.RUNTIME respectively.  If not, see
+   <http://www.gnu.org/licenses/>.  */
+
+/* This file implements timer routines for AMD GCN.  */
+
+#include "libgomp.h"
+
+/* According to AMD:
+    dGPU RTC is 27MHz
+    AGPU RTC is 100MHz
+   FIXME: DTRT on an APU.  */
+#define RTC_TICKS (1.0 / 27000000.0) /* 27MHz */
+
+double
+omp_get_wtime (void)
+{
+  uint64_t clock;
+  asm ("s_memrealtime %0\n\t"
+       "s_waitcnt 0" : "=r" (clock));
+  return clock * RTC_TICKS;
+}
+
+double
+omp_get_wtick (void)
+{
+  return RTC_TICKS;
+}
+
+ialias (omp_get_wtime)
+ialias (omp_get_wtick)
diff --git a/libgomp/configure.ac b/libgomp/configure.ac
index 0f9e821c007..725f3bfd285 100644
--- a/libgomp/configure.ac
+++ b/libgomp/configure.ac
@@ -176,7 +176,7 @@ case "$host" in
   *-*-rtems*)
     # RTEMS supports Pthreads, but the library is not available at GCC build time.
     ;;
-  nvptx*-*-*)
+  nvptx*-*-* | amdgcn*-*-*)
     # NVPTX does not support Pthreads, has its own code replacement.
     libgomp_use_pthreads=no
     # NVPTX is an accelerator-only target
diff --git a/libgomp/configure.tgt b/libgomp/configure.tgt
index c5ae9a9e39a..06ee115ece9 100644
--- a/libgomp/configure.tgt
+++ b/libgomp/configure.tgt
@@ -164,6 +164,10 @@ case "${target}" in
  fi
  ;;
 
+  amdgcn*-*-*)
+ config_path="gcn accel"
+ ;;
+
   *)
  ;;
 
diff --git a/libgomp/libgomp-plugin.h b/libgomp/libgomp-plugin.h
index de969e1ba45..037558c43f5 100644
--- a/libgomp/libgomp-plugin.h
+++ b/libgomp/libgomp-plugin.h
@@ -50,7 +50,8 @@ enum offload_target_type
   /* OFFLOAD_TARGET_TYPE_HOST_NONSHM = 3 removed.  */
   OFFLOAD_TARGET_TYPE_NVIDIA_PTX = 5,
   OFFLOAD_TARGET_TYPE_INTEL_MIC = 6,
-  OFFLOAD_TARGET_TYPE_HSA = 7
+  OFFLOAD_TARGET_TYPE_HSA = 7,
+  OFFLOAD_TARGET_TYPE_GCN = 8
 };
 
 /* Opaque type to represent plugin-dependent implementation of an
diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h
index 178eb600ccd..19e1241ee4c 100644
--- a/libgomp/libgomp.h
+++ b/libgomp/libgomp.h
@@ -692,6 +692,24 @@ static inline struct gomp_thread *gomp_thread (void)
   asm ("mov.u32 %0, %%tid.y;" : "=r" (tid));
   return nvptx_thrs + tid;
 }
+#elif defined __AMDGCN__
+static inline struct gomp_thread *gcn_thrs (void)
+{
+  /* The value is at the bottom of LDS.  */
+  struct gomp_thread * __lds *thrs = (struct gomp_thread * __lds *)4;
+  return *thrs;
+}
+static inline void set_gcn_thrs (struct gomp_thread *val)
+{
+  /* The value is at the bottom of LDS.  */
+  struct gomp_thread * __lds *thrs = (struct gomp_thread * __lds *)4;
+  *thrs = val;
+}
+static inline struct gomp_thread *gomp_thread (void)
+{
+  int tid = __builtin_gcn_dim_pos(1);
+  return gcn_thrs () + tid;
+}
 #elif defined HAVE_TLS || defined USE_EMUTLS
 extern __thread struct gomp_thread gomp_tls_data;
 static inline struct gomp_thread *gomp_thread (void)
diff --git a/libgomp/oacc-int.h b/libgomp/oacc-int.h
index 5ca9944601e..9dc6c8a5713 100644
--- a/libgomp/oacc-int.h
+++ b/libgomp/oacc-int.h
@@ -82,7 +82,14 @@ struct goacc_thread
   void *target_tls;
 };
 
-#if defined HAVE_TLS || defined USE_EMUTLS
+#ifdef __AMDGCN__
+static inline struct goacc_thread *
+goacc_thread (void)
+{
+  /* Unused in the offload libgomp for OpenACC: return a dummy value.  */
+  return 0;
+}
+#elif defined HAVE_TLS || defined USE_EMUTLS
 extern __thread struct goacc_thread *goacc_tls_data;
 static inline struct goacc_thread *
 goacc_thread (void)
diff --git a/libgomp/oacc-target.c b/libgomp/oacc-target.c
new file mode 100644
index 00000000000..f2e79899030
--- /dev/null
+++ b/libgomp/oacc-target.c
@@ -0,0 +1 @@
+/* Nothing needed here.  */
diff --git a/libgomp/openacc.f90 b/libgomp/openacc.f90
index bc205453f82..831a157e703 100644
--- a/libgomp/openacc.f90
+++ b/libgomp/openacc.f90
@@ -46,6 +46,7 @@ module openacc_kinds
   ! integer (acc_device_kind), parameter :: acc_device_host_nonshm = 3 removed.
   integer (acc_device_kind), parameter :: acc_device_not_host = 4
   integer (acc_device_kind), parameter :: acc_device_nvidia = 5
+  integer (acc_device_kind), parameter :: acc_device_gcn = 8
 
   public :: acc_handle_kind
 
diff --git a/libgomp/openacc.h b/libgomp/openacc.h
index 1bbe6c90e7f..42c861caabf 100644
--- a/libgomp/openacc.h
+++ b/libgomp/openacc.h
@@ -55,6 +55,7 @@ typedef enum acc_device_t {
   /* acc_device_host_nonshm = 3 removed.  */
   acc_device_not_host = 4,
   acc_device_nvidia = 5,
+  acc_device_gcn = 8,
   _ACC_device_hwm,
   /* Ensure enumeration is layout compatible with int.  */
   _ACC_highest = __INT_MAX__,
diff --git a/libgomp/team.c b/libgomp/team.c
index c422da3701d..b26caaaaec6 100644
--- a/libgomp/team.c
+++ b/libgomp/team.c
@@ -239,6 +239,9 @@ gomp_free_pool_helper (void *thread_pool)
   pthread_exit (NULL);
 #elif defined(__nvptx__)
   asm ("exit;");
+#elif defined(__AMDGCN__)
+  asm ("s_dcache_wb\n\t"
+       "s_endpgm");
 #else
 #error gomp_free_pool_helper must terminate the thread
 #endif
Reply | Threaded
Open this post in threaded view
|

[PATCH 5/7 libgomp,amdgcn] Optimize GCN OpenMP malloc performance

Andrew Stubbs-4
In reply to this post by Andrew Stubbs-4
This patch implements a malloc optimization to improve the startup and
shutdown overhead for each OpenMP team.

New malloc functions are created, "team_malloc" and "team_free", that
take memory from a per-team memory arena provided by the plugin, rather
than the shared heap space, which is slow, and gets worse the more teams
are trying to allocate at once.

These new functions are used both in the gcn/team.c file and in selected
places elsewhere in libgomp.  Arena-space is limited (and larger sizes
have greater overhead at launch time) so this should not be a global
search and replace.

Dummy pass-through definitions are provided for other targets.

OK to commit?

Thanks

Andrew


2019-11-12  Andrew Stubbs  <[hidden email]>

        libgomp/
        * config/gcn/team.c (gomp_gcn_enter_kernel): Set up the team arena
        and use team_malloc variants.
        (gomp_gcn_exit_kernel): Use team_free.
        * libgomp.h (TEAM_ARENA_SIZE): Define.
        (TEAM_ARENA_FREE): Define.
        (TEAM_ARENA_END): Define.
        (team_malloc): New function.
        (team_malloc_cleared): New function.
        (team_free): New function.
        * team.c (gomp_new_team): Use team_malloc.
        (free_team): Use team_free.
        (gomp_free_thread): Use team_free.
        (gomp_pause_host): Use team_free.
        * work.c (gomp_init_work_share): Use team_malloc.
        (gomp_fini_work_share): Use team_free.
---
 libgomp/config/gcn/team.c | 18 ++++++++++---
 libgomp/libgomp.h         | 56 +++++++++++++++++++++++++++++++++++++++
 libgomp/team.c            | 12 ++++-----
 libgomp/work.c            |  4 +--
 4 files changed, 78 insertions(+), 12 deletions(-)


diff --git a/libgomp/config/gcn/team.c b/libgomp/config/gcn/team.c
index c566482bda2..063571fc751 100644
--- a/libgomp/config/gcn/team.c
+++ b/libgomp/config/gcn/team.c
@@ -57,16 +57,26 @@ gomp_gcn_enter_kernel (void)
       /* Starting additional threads is not supported.  */
       gomp_global_icv.dyn_var = true;
 
+      /* Initialize the team arena for optimized memory allocation.
+         The arena has been allocated on the host side, and the address
+         passed in via the kernargs.  Each team takes a small slice of it.  */
+      register void **kernargs asm("s8");
+      void *team_arena = (kernargs[4] + TEAM_ARENA_SIZE*teamid);
+      void * __lds *arena_free = (void * __lds *)TEAM_ARENA_FREE;
+      void * __lds *arena_end = (void * __lds *)TEAM_ARENA_END;
+      *arena_free = team_arena;
+      *arena_end = team_arena + TEAM_ARENA_SIZE;
+
       /* Allocate and initialize the team-local-storage data.  */
-      struct gomp_thread *thrs = gomp_malloc_cleared (sizeof (*thrs)
+      struct gomp_thread *thrs = team_malloc_cleared (sizeof (*thrs)
       * numthreads);
       set_gcn_thrs (thrs);
 
       /* Allocate and initailize a pool of threads in the team.
          The threads are already running, of course, we just need to manage
          the communication between them.  */
-      struct gomp_thread_pool *pool = gomp_malloc (sizeof (*pool));
-      pool->threads = gomp_malloc (sizeof (void *) * numthreads);
+      struct gomp_thread_pool *pool = team_malloc (sizeof (*pool));
+      pool->threads = team_malloc (sizeof (void *) * numthreads);
       for (int tid = 0; tid < numthreads; tid++)
  pool->threads[tid] = &thrs[tid];
       pool->threads_size = numthreads;
@@ -91,7 +101,7 @@ void
 gomp_gcn_exit_kernel (void)
 {
   gomp_free_thread (gcn_thrs ());
-  free (gcn_thrs ());
+  team_free (gcn_thrs ());
 }
 
 /* This function contains the idle loop in which a thread waits
diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h
index 19e1241ee4c..659aeb95ffe 100644
--- a/libgomp/libgomp.h
+++ b/libgomp/libgomp.h
@@ -106,6 +106,62 @@ extern void gomp_aligned_free (void *);
    GCC's builtin alloca().  */
 #define gomp_alloca(x)  __builtin_alloca(x)
 
+/* Optimized allocators for team-specific data that will die with the team.  */
+
+#ifdef __AMDGCN__
+/* The arena is initialized in config/gcn/team.c.  */
+#define TEAM_ARENA_SIZE 64*1024  /* Must match the value in plugin-gcn.c.  */
+#define TEAM_ARENA_FREE 16  /* LDS offset of free pointer.  */
+#define TEAM_ARENA_END  24  /* LDS offset of end pointer.  */
+
+static inline void * __attribute__((malloc))
+team_malloc (size_t size)
+{
+  /* 4-byte align the size.  */
+  size = (size + 3) & ~3;
+
+  /* Allocate directly from the arena.
+     The compiler does not support DS atomics, yet. */
+  void *result;
+  asm ("ds_add_rtn_u64 %0, %1, %2\n\ts_waitcnt 0"
+       : "=v"(result) : "v"(TEAM_ARENA_FREE), "v"(size), "e"(1L) : "memory");
+
+  /* Handle OOM.  */
+  if (result + size > *(void * __lds *)TEAM_ARENA_END)
+    {
+      const char msg[] = "GCN team arena exhausted\n";
+      write (2, msg, sizeof(msg)-1);
+      /* It's better to continue with reeduced performance than abort.
+         Beware that this won't get freed, which might cause more problems.  */
+      result = gomp_malloc (size);
+    }
+  return result;
+}
+
+static inline void * __attribute__((malloc)) __attribute__((optimize("-O3")))
+team_malloc_cleared (size_t size)
+{
+  char *result = team_malloc (size);
+
+  /* Clear the allocated memory.
+     This should vectorize.  The allocation has been rounded up to the next
+     4-byte boundary, so this is safe.  */
+  for (int i = 0; i<size; i+=4)
+    *(int*)(result+i) = 0;
+  return result;
+}
+
+static inline void
+team_free (void *ptr)
+{
+  /* The whole arena is freed when the kernel exits.  */
+}
+#else
+#define team_malloc(...) gomp_malloc (__VA_ARGS__)
+#define team_malloc_cleared(...) gomp_malloc_cleared (__VA_ARGS__)
+#define team_free(...) free (__VA_ARGS__)
+#endif
+
 /* error.c */
 
 extern void gomp_vdebug (int, const char *, va_list);
diff --git a/libgomp/team.c b/libgomp/team.c
index b26caaaaec6..cdfb9ba6c98 100644
--- a/libgomp/team.c
+++ b/libgomp/team.c
@@ -171,7 +171,7 @@ gomp_new_team (unsigned nthreads)
     {
       size_t extra = sizeof (team->ordered_release[0])
      + sizeof (team->implicit_task[0]);
-      team = gomp_malloc (sizeof (*team) + nthreads * extra);
+      team = team_malloc (sizeof (*team) + nthreads * extra);
 
 #ifndef HAVE_SYNC_BUILTINS
       gomp_mutex_init (&team->work_share_list_free_lock);
@@ -221,7 +221,7 @@ free_team (struct gomp_team *team)
   gomp_barrier_destroy (&team->barrier);
   gomp_mutex_destroy (&team->task_lock);
   priority_queue_free (&team->task_queue);
-  free (team);
+  team_free (team);
 }
 
 static void
@@ -285,8 +285,8 @@ gomp_free_thread (void *arg __attribute__((unused)))
       if (pool->last_team)
  free_team (pool->last_team);
 #ifndef __nvptx__
-      free (pool->threads);
-      free (pool);
+      team_free (pool->threads);
+      team_free (pool);
 #endif
       thr->thread_pool = NULL;
     }
@@ -1082,8 +1082,8 @@ gomp_pause_host (void)
       if (pool->last_team)
  free_team (pool->last_team);
 #ifndef __nvptx__
-      free (pool->threads);
-      free (pool);
+      team_free (pool->threads);
+      team_free (pool);
 #endif
       thr->thread_pool = NULL;
     }
diff --git a/libgomp/work.c b/libgomp/work.c
index a589b8b5231..28bb0c11255 100644
--- a/libgomp/work.c
+++ b/libgomp/work.c
@@ -120,7 +120,7 @@ gomp_init_work_share (struct gomp_work_share *ws, size_t ordered,
       else
  ordered = nthreads * sizeof (*ws->ordered_team_ids);
       if (ordered > INLINE_ORDERED_TEAM_IDS_SIZE)
- ws->ordered_team_ids = gomp_malloc (ordered);
+ ws->ordered_team_ids = team_malloc (ordered);
       else
  ws->ordered_team_ids = ws->inline_ordered_team_ids;
       memset (ws->ordered_team_ids, '\0', ordered);
@@ -142,7 +142,7 @@ gomp_fini_work_share (struct gomp_work_share *ws)
 {
   gomp_mutex_destroy (&ws->lock);
   if (ws->ordered_team_ids != ws->inline_ordered_team_ids)
-    free (ws->ordered_team_ids);
+    team_free (ws->ordered_team_ids);
   gomp_ptrlock_destroy (&ws->next_ws);
 }
 
Reply | Threaded
Open this post in threaded view
|

[PATCH 6/7 amdgcn] Use a single worker for OpenACC on AMD GCN

Andrew Stubbs-4
In reply to this post by Andrew Stubbs-4
This patch prevents the compiler using multiple workers in a gang.  This
should be reverted when worker support is committed.

I will commit this with the reset of the series.

Andrew


2019-11-12  Andrew Stubbs  <[hidden email]>
            Julian Brown  <[hidden email]>

        gcc/
        * config/gcn/gcn.c (gcn_goacc_validate_dims): Ensure
        flag_worker_partitioning is not set.
        (TARGET_GOACC_WORKER_PARTITIONING): Remove target hook definition.
        * config/gcn/gcn.opt (macc-experimental-workers): Default to off.
---
 gcc/config/gcn/gcn.c   | 4 ++--
 gcc/config/gcn/gcn.opt | 2 +-
 2 files changed, 3 insertions(+), 3 deletions(-)


diff --git a/gcc/config/gcn/gcn.c b/gcc/config/gcn/gcn.c
index cdd24277cf6..1a69737f693 100644
--- a/gcc/config/gcn/gcn.c
+++ b/gcc/config/gcn/gcn.c
@@ -4695,6 +4695,8 @@ gcn_goacc_validate_dims (tree decl, int dims[], int fn_level,
   /* FIXME: remove -facc-experimental-workers when they're ready.  */
   int max_workers = flag_worker_partitioning ? 16 : 1;
 
+  gcc_assert (!flag_worker_partitioning);
+
   /* The vector size must appear to be 64, to the user, unless this is a
      SEQ routine.  The real, internal value is always 1, which means use
      autovectorization, but the user should not see that.  */
@@ -6073,8 +6075,6 @@ print_operand (FILE *file, rtx x, int code)
 #define TARGET_GOACC_REDUCTION gcn_goacc_reduction
 #undef  TARGET_GOACC_VALIDATE_DIMS
 #define TARGET_GOACC_VALIDATE_DIMS gcn_goacc_validate_dims
-#undef  TARGET_GOACC_WORKER_PARTITIONING
-#define TARGET_GOACC_WORKER_PARTITIONING true
 #undef  TARGET_HARD_REGNO_MODE_OK
 #define TARGET_HARD_REGNO_MODE_OK gcn_hard_regno_mode_ok
 #undef  TARGET_HARD_REGNO_NREGS
diff --git a/gcc/config/gcn/gcn.opt b/gcc/config/gcn/gcn.opt
index bdc878f35ad..402deb625bd 100644
--- a/gcc/config/gcn/gcn.opt
+++ b/gcc/config/gcn/gcn.opt
@@ -65,7 +65,7 @@ Target Report RejectNegative Var(flag_bypass_init_error)
 bool flag_worker_partitioning = false
 
 macc-experimental-workers
-Target Report Var(flag_worker_partitioning) Init(1)
+Target Report Var(flag_worker_partitioning) Init(0)
 
 int stack_size_opt = -1
 
Reply | Threaded
Open this post in threaded view
|

[PATCH 7/7 libgomp,amdgcn] GCN Libgomp Plugin

Andrew Stubbs-4
In reply to this post by Andrew Stubbs-4
This patch contributes the GCN libgomp plugin, with the various
configure and make bits to go with it.

This implementation is a much-cleaned-up version of the one present on the
openacc-gcc-9-branch.

OK to commit?

Thanks

Andrew


2019-11-12  Andrew Stubbs  <[hidden email]>

        libgomp/
        * plugin/Makefrag.am: Add amdgcn plugin support.
        * plugin/configfrag.ac: Likewise.
        * plugin/plugin-gcn.c: New file.
        * configure: Regenerate.
---
 libgomp/plugin/Makefrag.am   |   14 +
 libgomp/plugin/configfrag.ac |   35 +
 libgomp/plugin/plugin-gcn.c  | 3985 ++++++++++++++++++++++++++++++++++
 3 files changed, 4034 insertions(+)
 create mode 100644 libgomp/plugin/plugin-gcn.c


diff --git a/libgomp/plugin/Makefrag.am b/libgomp/plugin/Makefrag.am
index 168ef59de41..45ed043e333 100644
--- a/libgomp/plugin/Makefrag.am
+++ b/libgomp/plugin/Makefrag.am
@@ -52,3 +52,17 @@ libgomp_plugin_hsa_la_LDFLAGS += $(PLUGIN_HSA_LDFLAGS)
 libgomp_plugin_hsa_la_LIBADD = libgomp.la $(PLUGIN_HSA_LIBS)
 libgomp_plugin_hsa_la_LIBTOOLFLAGS = --tag=disable-static
 endif
+
+if PLUGIN_GCN
+# AMD GCN plugin
+libgomp_plugin_gcn_version_info = -version-info $(libtool_VERSION)
+toolexeclib_LTLIBRARIES += libgomp-plugin-gcn.la
+libgomp_plugin_gcn_la_SOURCES = plugin/plugin-gcn.c
+libgomp_plugin_gcn_la_CPPFLAGS = $(AM_CPPFLAGS) $(PLUGIN_GCN_CPPFLAGS) \
+ -D_GNU_SOURCE
+libgomp_plugin_gcn_la_LDFLAGS = $(libgomp_plugin_gcn_version_info) \
+ $(lt_host_flags)
+libgomp_plugin_gcn_la_LDFLAGS += $(PLUGIN_GCN_LDFLAGS)
+libgomp_plugin_gcn_la_LIBADD = libgomp.la $(PLUGIN_GCN_LIBS)
+libgomp_plugin_gcn_la_LIBTOOLFLAGS = --tag=disable-static
+endif
diff --git a/libgomp/plugin/configfrag.ac b/libgomp/plugin/configfrag.ac
index 9718ac752e2..424ec6c96b2 100644
--- a/libgomp/plugin/configfrag.ac
+++ b/libgomp/plugin/configfrag.ac
@@ -137,6 +137,15 @@ AC_SUBST(PLUGIN_HSA_CPPFLAGS)
 AC_SUBST(PLUGIN_HSA_LDFLAGS)
 AC_SUBST(PLUGIN_HSA_LIBS)
 
+PLUGIN_GCN=0
+PLUGIN_GCN_CPPFLAGS=
+PLUGIN_GCN_LDFLAGS=
+PLUGIN_GCN_LIBS=
+AC_SUBST(PLUGIN_GCN)
+AC_SUBST(PLUGIN_GCN_CPPFLAGS)
+AC_SUBST(PLUGIN_GCN_LDFLAGS)
+AC_SUBST(PLUGIN_GCN_LIBS)
+
 # Parse '--enable-offload-targets', figure out the corresponding libgomp
 # plugins, and configure to find the corresponding offload compilers.
 # 'offload_plugins' and 'offload_targets' will be populated in the same order.
@@ -237,6 +246,29 @@ if test x"$enable_offload_targets" != x; then
             ;;
         esac
         ;;
+
+      amdgcn*)
+ case "${target}" in
+  x86_64-*-*)
+    case " ${CC} ${CFLAGS} " in
+      *" -m32 "*)
+ PLUGIN_GCN=0
+ ;;
+      *)
+ tgt_plugin=gcn
+ PLUGIN_GCN=$tgt
+ PLUGIN_GCN_CPPFLAGS=$HSA_RUNTIME_CPPFLAGS
+ PLUGIN_GCN_LDFLAGS="$HSA_RUNTIME_LDFLAGS"
+ PLUGIN_GCN_LIBS="-ldl"
+ PLUGIN_GCN=1
+ ;;
+      esac
+    ;;
+  *-*-*)
+    PLUGIN_GCN=0
+     ;;
+ esac
+ ;;
       *)
  AC_MSG_ERROR([unknown offload target specified])
  ;;
@@ -275,6 +307,9 @@ AC_DEFINE_UNQUOTED([PLUGIN_NVPTX_DYNAMIC], [$PLUGIN_NVPTX_DYNAMIC],
 AM_CONDITIONAL([PLUGIN_HSA], [test $PLUGIN_HSA = 1])
 AC_DEFINE_UNQUOTED([PLUGIN_HSA], [$PLUGIN_HSA],
   [Define to 1 if the HSA plugin is built, 0 if not.])
+AM_CONDITIONAL([PLUGIN_GCN], [test $PLUGIN_GCN = 1])
+AC_DEFINE_UNQUOTED([PLUGIN_GCN], [$PLUGIN_GCN],
+  [Define to 1 if the GCN plugin is built, 0 if not.])
 
 if test "$HSA_RUNTIME_LIB" != ""; then
   HSA_RUNTIME_LIB="$HSA_RUNTIME_LIB/"
diff --git a/libgomp/plugin/plugin-gcn.c b/libgomp/plugin/plugin-gcn.c
new file mode 100644
index 00000000000..583916759a5
--- /dev/null
+++ b/libgomp/plugin/plugin-gcn.c
@@ -0,0 +1,3985 @@
+/* Plugin for AMD GCN execution.
+
+   Copyright (C) 2013-2019 Free Software Foundation, Inc.
+
+   Contributed by Mentor Embedded
+
+   This file is part of the GNU Offloading and Multi Processing Library
+   (libgomp).
+
+   Libgomp is free software; you can redistribute it and/or modify it
+   under the terms of the GNU General Public License as published by
+   the Free Software Foundation; either version 3, or (at your option)
+   any later version.
+
+   Libgomp is distributed in the hope that it will be useful, but WITHOUT ANY
+   WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS
+   FOR A PARTICULAR PURPOSE.  See the GNU General Public License for
+   more details.
+
+   Under Section 7 of GPL version 3, you are granted additional
+   permissions described in the GCC Runtime Library Exception, version
+   3.1, as published by the Free Software Foundation.
+
+   You should have received a copy of the GNU General Public License and
+   a copy of the GCC Runtime Library Exception along with this program;
+   see the files COPYING3 and COPYING.RUNTIME respectively.  If not, see
+   <http://www.gnu.org/licenses/>.  */
+
+/* {{{ Includes and defines  */
+
+#include "config.h"
+#include <stdio.h>
+#include <stdlib.h>
+#include <string.h>
+#include <pthread.h>
+#include <inttypes.h>
+#include <stdbool.h>
+#include <limits.h>
+#include <hsa.h>
+#include <dlfcn.h>
+#include <signal.h>
+#include "libgomp-plugin.h"
+#include "gomp-constants.h"
+#include <elf.h>
+#include "oacc-plugin.h"
+#include "oacc-int.h"
+#include <assert.h>
+
+/* Additional definitions not in HSA 1.1.
+   FIXME: this needs to be updated in hsa.h for upstream, but the only source
+          right now is the ROCr source which may cause license issues.  */
+#define HSA_AMD_AGENT_INFO_COMPUTE_UNIT_COUNT 0xA002
+
+/* These probably won't be in elf.h for a while.  */
+#define R_AMDGPU_NONE 0
+#define R_AMDGPU_ABS32_LO 1 /* (S + A) & 0xFFFFFFFF  */
+#define R_AMDGPU_ABS32_HI 2 /* (S + A) >> 32  */
+#define R_AMDGPU_ABS64 3 /* S + A  */
+#define R_AMDGPU_REL32 4 /* S + A - P  */
+#define R_AMDGPU_REL64 5 /* S + A - P  */
+#define R_AMDGPU_ABS32 6 /* S + A  */
+#define R_AMDGPU_GOTPCREL 7 /* G + GOT + A - P  */
+#define R_AMDGPU_GOTPCREL32_LO 8 /* (G + GOT + A - P) & 0xFFFFFFFF  */
+#define R_AMDGPU_GOTPCREL32_HI 9 /* (G + GOT + A - P) >> 32  */
+#define R_AMDGPU_REL32_LO 10 /* (S + A - P) & 0xFFFFFFFF  */
+#define R_AMDGPU_REL32_HI 11 /* (S + A - P) >> 32  */
+#define reserved 12
+#define R_AMDGPU_RELATIVE64 13 /* B + A  */
+
+/* GCN specific definitions for asynchronous queues.  */
+
+#define ASYNC_QUEUE_SIZE 64
+#define DRAIN_QUEUE_SYNCHRONOUS_P false
+#define DEBUG_QUEUES 0
+#define DEBUG_THREAD_SLEEP 0
+#define DEBUG_THREAD_SIGNAL 0
+
+/* Defaults.  */
+#define DEFAULT_GCN_HEAP_SIZE (100*1024*1024)  /* 100MB.  */
+
+/* Secure getenv() which returns NULL if running as SUID/SGID.  */
+#ifndef HAVE_SECURE_GETENV
+#ifdef HAVE___SECURE_GETENV
+#define secure_getenv __secure_getenv
+#elif defined (HAVE_UNISTD_H) && defined(HAVE_GETUID) && defined(HAVE_GETEUID) \
+  && defined(HAVE_GETGID) && defined(HAVE_GETEGID)
+
+#include <unistd.h>
+
+/* Implementation of secure_getenv() for targets where it is not provided but
+   we have at least means to test real and effective IDs. */
+
+static char *
+secure_getenv (const char *name)
+{
+  if ((getuid () == geteuid ()) && (getgid () == getegid ()))
+    return getenv (name);
+  else
+    return NULL;
+}
+
+#else
+#define secure_getenv getenv
+#endif
+#endif
+
+/* }}}  */
+/* {{{ Types  */
+
+/* GCN-specific implmentation of the GOMP_PLUGIN_acc_thread data.  */
+
+struct gcn_thread
+{
+  /* The thread number from the async clause, or GOMP_ASYNC_SYNC.  */
+  int async;
+};
+
+/* As an HSA runtime is dlopened, following structure defines function
+   pointers utilized by the HSA plug-in.  */
+
+struct hsa_runtime_fn_info
+{
+  /* HSA runtime.  */
+  hsa_status_t (*hsa_status_string_fn) (hsa_status_t status,
+ const char **status_string);
+  hsa_status_t (*hsa_system_get_info_fn) (hsa_system_info_t attribute,
+  void *value);
+  hsa_status_t (*hsa_agent_get_info_fn) (hsa_agent_t agent,
+ hsa_agent_info_t attribute,
+ void *value);
+  hsa_status_t (*hsa_isa_get_info_fn)(hsa_isa_t isa,
+      hsa_isa_info_t attribute,
+      uint32_t index,
+      void *value);
+  hsa_status_t (*hsa_init_fn) (void);
+  hsa_status_t (*hsa_iterate_agents_fn)
+    (hsa_status_t (*callback)(hsa_agent_t agent, void *data), void *data);
+  hsa_status_t (*hsa_region_get_info_fn) (hsa_region_t region,
+  hsa_region_info_t attribute,
+  void *value);
+  hsa_status_t (*hsa_queue_create_fn)
+    (hsa_agent_t agent, uint32_t size, hsa_queue_type_t type,
+     void (*callback)(hsa_status_t status, hsa_queue_t *source, void *data),
+     void *data, uint32_t private_segment_size,
+     uint32_t group_segment_size, hsa_queue_t **queue);
+  hsa_status_t (*hsa_agent_iterate_regions_fn)
+    (hsa_agent_t agent,
+     hsa_status_t (*callback)(hsa_region_t region, void *data), void *data);
+  hsa_status_t (*hsa_executable_destroy_fn) (hsa_executable_t executable);
+  hsa_status_t (*hsa_executable_create_fn)
+    (hsa_profile_t profile, hsa_executable_state_t executable_state,
+     const char *options, hsa_executable_t *executable);
+  hsa_status_t (*hsa_executable_global_variable_define_fn)
+    (hsa_executable_t executable, const char *variable_name, void *address);
+  hsa_status_t (*hsa_executable_load_code_object_fn)
+    (hsa_executable_t executable, hsa_agent_t agent,
+     hsa_code_object_t code_object, const char *options);
+  hsa_status_t (*hsa_executable_freeze_fn)(hsa_executable_t executable,
+   const char *options);
+  hsa_status_t (*hsa_signal_create_fn) (hsa_signal_value_t initial_value,
+ uint32_t num_consumers,
+ const hsa_agent_t *consumers,
+ hsa_signal_t *signal);
+  hsa_status_t (*hsa_memory_allocate_fn) (hsa_region_t region, size_t size,
+  void **ptr);
+  hsa_status_t (*hsa_memory_assign_agent_fn) (void *ptr, hsa_agent_t agent,
+      hsa_access_permission_t access);
+  hsa_status_t (*hsa_memory_copy_fn)(void *dst, const void *src, size_t size);
+  hsa_status_t (*hsa_memory_free_fn) (void *ptr);
+  hsa_status_t (*hsa_signal_destroy_fn) (hsa_signal_t signal);
+  hsa_status_t (*hsa_executable_get_symbol_fn)
+    (hsa_executable_t executable, const char *module_name,
+     const char *symbol_name, hsa_agent_t agent, int32_t call_convention,
+     hsa_executable_symbol_t *symbol);
+  hsa_status_t (*hsa_executable_symbol_get_info_fn)
+    (hsa_executable_symbol_t executable_symbol,
+     hsa_executable_symbol_info_t attribute, void *value);
+  hsa_status_t (*hsa_executable_iterate_symbols_fn)
+    (hsa_executable_t executable,
+     hsa_status_t (*callback)(hsa_executable_t executable,
+      hsa_executable_symbol_t symbol, void *data),
+     void *data);
+  uint64_t (*hsa_queue_add_write_index_release_fn) (const hsa_queue_t *queue,
+    uint64_t value);
+  uint64_t (*hsa_queue_load_read_index_acquire_fn) (const hsa_queue_t *queue);
+  void (*hsa_signal_store_relaxed_fn) (hsa_signal_t signal,
+       hsa_signal_value_t value);
+  void (*hsa_signal_store_release_fn) (hsa_signal_t signal,
+       hsa_signal_value_t value);
+  hsa_signal_value_t (*hsa_signal_wait_acquire_fn)
+    (hsa_signal_t signal, hsa_signal_condition_t condition,
+     hsa_signal_value_t compare_value, uint64_t timeout_hint,
+     hsa_wait_state_t wait_state_hint);
+  hsa_signal_value_t (*hsa_signal_load_acquire_fn) (hsa_signal_t signal);
+  hsa_status_t (*hsa_queue_destroy_fn) (hsa_queue_t *queue);
+
+  hsa_status_t (*hsa_code_object_deserialize_fn)
+    (void *serialized_code_object, size_t serialized_code_object_size,
+     const char *options, hsa_code_object_t *code_object);
+};
+
+/* Structure describing the run-time and grid properties of an HSA kernel
+   lauch.  This needs to match the format passed to GOMP_OFFLOAD_run.  */
+
+struct GOMP_kernel_launch_attributes
+{
+  /* Number of dimensions the workload has.  Maximum number is 3.  */
+  uint32_t ndim;
+  /* Size of the grid in the three respective dimensions.  */
+  uint32_t gdims[3];
+  /* Size of work-groups in the respective dimensions.  */
+  uint32_t wdims[3];
+};
+
+/* Collection of information needed for a dispatch of a kernel from a
+   kernel.  */
+
+struct kernel_dispatch
+{
+  struct agent_info *agent;
+  /* Pointer to a command queue associated with a kernel dispatch agent.  */
+  void *queue;
+  /* Pointer to a memory space used for kernel arguments passing.  */
+  void *kernarg_address;
+  /* Kernel object.  */
+  uint64_t object;
+  /* Synchronization signal used for dispatch synchronization.  */
+  uint64_t signal;
+  /* Private segment size.  */
+  uint32_t private_segment_size;
+  /* Group segment size.  */
+  uint32_t group_segment_size;
+};
+
+/* Structure of the kernargs segment, supporting console output.
+
+   This needs to match the definitions in Newlib, and the expectations
+   in libgomp target code.  */
+
+struct kernargs {
+  /* Leave space for the real kernel arguments.
+     OpenACC and OpenMP only use one pointer.  */
+  int64_t dummy1;
+  int64_t dummy2;
+
+  /* A pointer to struct output, below, for console output data.  */
+  int64_t out_ptr;
+
+  /* A pointer to struct heap, below.  */
+  int64_t heap_ptr;
+
+  /* A pointer to an ephemeral memory arena.
+    Only needed for OpenMP.  */
+  int64_t arena_ptr;
+
+  /* Output data.  */
+  struct output {
+    int return_value;
+    unsigned int next_output;
+    struct printf_data {
+      int written;
+      char msg[128];
+      int type;
+      union {
+ int64_t ivalue;
+ double dvalue;
+ char text[128];
+      };
+    } queue[1024];
+    unsigned int consumed;
+  } output_data;
+};
+
+/* A queue entry for a future asynchronous launch.  */
+
+struct kernel_launch
+{
+  struct kernel_info *kernel;
+  void *vars;
+  struct GOMP_kernel_launch_attributes kla;
+};
+
+/* A queue entry for a future callback.  */
+
+struct callback
+{
+  void (*fn)(void *);
+  void *data;
+};
+
+/* A data struct for the copy_data callback.  */
+
+struct copy_data
+{
+  void *dst;
+  const void *src;
+  size_t len;
+  bool free_src;
+  struct goacc_asyncqueue *aq;
+};
+
+/* A queue entry for a placeholder.  These correspond to a wait event.  */
+
+struct placeholder
+{
+  int executed;
+  pthread_cond_t cond;
+  pthread_mutex_t mutex;
+};
+
+/* A queue entry for a wait directive.  */
+
+struct asyncwait_info
+{
+  struct placeholder *placeholderp;
+};
+
+/* Encode the type of an entry in an async queue.  */
+
+enum entry_type
+{
+  KERNEL_LAUNCH,
+  CALLBACK,
+  ASYNC_WAIT,
+  ASYNC_PLACEHOLDER
+};
+
+/* An entry in an async queue.  */
+
+struct queue_entry
+{
+  enum entry_type type;
+  union {
+    struct kernel_launch launch;
+    struct callback callback;
+    struct asyncwait_info asyncwait;
+    struct placeholder placeholder;
+  } u;
+};
+
+/* An async queue header.
+
+   OpenMP may create one of these.
+   OpenACC may create many.  */
+
+struct goacc_asyncqueue
+{
+  struct agent_info *agent;
+  hsa_queue_t *hsa_queue;
+
+  pthread_t thread_drain_queue;
+  pthread_mutex_t mutex;
+  pthread_cond_t queue_cond_in;
+  pthread_cond_t queue_cond_out;
+  struct queue_entry queue[ASYNC_QUEUE_SIZE];
+  int queue_first;
+  int queue_n;
+  int drain_queue_stop;
+
+  int id;
+  struct goacc_asyncqueue *prev;
+  struct goacc_asyncqueue *next;
+};
+
+/* Mkoffload uses this structure to describe a kernel.
+
+   OpenMP kernel dimensions are passed at runtime.
+   OpenACC kernel dimensions are passed at compile time, here.  */
+
+struct hsa_kernel_description
+{
+  const char *name;
+  int oacc_dims[3];  /* Only present for GCN kernels.  */
+};
+
+/* Mkoffload uses this structure to describe an offload variable.  */
+
+struct global_var_info
+{
+  const char *name;
+  void *address;
+};
+
+/* Mkoffload uses this structure to describe all the kernels in a
+   loadable module.  These are passed the libgomp via static constructors.  */
+
+struct gcn_image_desc
+{
+  struct gcn_image {
+    size_t size;
+    void *image;
+  } *gcn_image;
+  const unsigned kernel_count;
+  struct hsa_kernel_description *kernel_infos;
+  const unsigned global_variable_count;
+  struct global_var_info *global_variables;
+};
+
+/* Description of an HSA GPU agent (device) and the program associated with
+   it.  */
+
+struct agent_info
+{
+  /* The HSA ID of the agent.  Assigned when hsa_context is initialized.  */
+  hsa_agent_t id;
+  /* The user-visible device number.  */
+  int device_id;
+  /* Whether the agent has been initialized.  The fields below are usable only
+     if it has been.  */
+  bool initialized;
+  /* Precomuted check for problem architectures.  */
+  bool gfx900_p;
+
+  /* Command queues of the agent.  */
+  hsa_queue_t *sync_queue;
+  struct goacc_asyncqueue *async_queues, *omp_async_queue;
+  pthread_mutex_t async_queues_mutex;
+
+  /* The HSA memory region from which to allocate kernel arguments.  */
+  hsa_region_t kernarg_region;
+
+  /* The HSA memory region from which to allocate device data.  */
+  hsa_region_t data_region;
+
+  /* Allocated team arenas.  */
+  struct team_arena_list *team_arena_list;
+  pthread_mutex_t team_arena_write_lock;
+
+  /* Read-write lock that protects kernels which are running or about to be run
+     from interference with loading and unloading of images.  Needs to be
+     locked for reading while a kernel is being run, and for writing if the
+     list of modules is manipulated (and thus the HSA program invalidated).  */
+  pthread_rwlock_t module_rwlock;
+
+  /* The module associated with this kernel.  */
+  struct module_info *module;
+
+  /* Mutex enforcing that only one thread will finalize the HSA program.  A
+     thread should have locked agent->module_rwlock for reading before
+     acquiring it.  */
+  pthread_mutex_t prog_mutex;
+  /* Flag whether the HSA program that consists of all the modules has been
+     finalized.  */
+  bool prog_finalized;
+  /* HSA executable - the finalized program that is used to locate kernels.  */
+  hsa_executable_t executable;
+};
+
+/* Information required to identify, finalize and run any given kernel.  */
+
+enum offload_kind {KIND_UNKNOWN, KIND_OPENMP, KIND_OPENACC};
+
+struct kernel_info
+{
+  /* Name of the kernel, required to locate it within the GCN object-code
+     module.  */
+  const char *name;
+  /* The specific agent the kernel has been or will be finalized for and run
+     on.  */
+  struct agent_info *agent;
+  /* The specific module where the kernel takes place.  */
+  struct module_info *module;
+  /* Mutex enforcing that at most once thread ever initializes a kernel for
+     use.  A thread should have locked agent->module_rwlock for reading before
+     acquiring it.  */
+  pthread_mutex_t init_mutex;
+  /* Flag indicating whether the kernel has been initialized and all fields
+     below it contain valid data.  */
+  bool initialized;
+  /* Flag indicating that the kernel has a problem that blocks an execution.  */
+  bool initialization_failed;
+  /* The object to be put into the dispatch queue.  */
+  uint64_t object;
+  /* Required size of kernel arguments.  */
+  uint32_t kernarg_segment_size;
+  /* Required size of group segment.  */
+  uint32_t group_segment_size;
+  /* Required size of private segment.  */
+  uint32_t private_segment_size;
+  /* Set up for OpenMP or OpenACC?  */
+  enum offload_kind kind;
+};
+
+/* Information about a particular GCN module, its image and kernels.  */
+
+struct module_info
+{
+  /* The description with which the program has registered the image.  */
+  struct gcn_image_desc *image_desc;
+  /* GCN heap allocation.  */
+  struct heap *heap;
+  /* Physical boundaries of the loaded module.  */
+  Elf64_Addr phys_address_start;
+  Elf64_Addr phys_address_end;
+
+  bool constructors_run_p;
+  struct kernel_info *init_array_func, *fini_array_func;
+
+  /* Number of kernels in this module.  */
+  int kernel_count;
+  /* An array of kernel_info structures describing each kernel in this
+     module.  */
+  struct kernel_info kernels[];
+};
+
+/* A linked list of memory arenas allocated on the device.
+   These are only used by OpenMP, as a means to optimize per-team malloc.  */
+
+struct team_arena_list
+{
+  struct team_arena_list *next;
+
+  /* The number of teams determines the size of the allocation.  */
+  int num_teams;
+  /* The device address of the arena itself.  */
+  void *arena;
+  /* A flag to prevent two asynchronous kernels trying to use the same arena.
+     The mutex is locked until the kernel exits.  */
+  pthread_mutex_t in_use;
+};
+
+/* Information about the whole HSA environment and all of its agents.  */
+
+struct hsa_context_info
+{
+  /* Whether the structure has been initialized.  */
+  bool initialized;
+  /* Number of usable GPU HSA agents in the system.  */
+  int agent_count;
+  /* Array of agent_info structures describing the individual HSA agents.  */
+  struct agent_info *agents;
+};
+
+/* Format of the on-device heap.
+
+   This must match the definition in Newlib and gcn-run.  */
+
+struct heap {
+  int64_t size;
+  char data[0];
+};
+
+/* }}}  */
+/* {{{ Global variables  */
+
+/* Information about the whole HSA environment and all of its agents.  */
+
+static struct hsa_context_info hsa_context;
+
+/* HSA runtime functions that are initialized in init_hsa_context.  */
+
+static struct hsa_runtime_fn_info hsa_fns;
+
+/* Heap space, allocated target-side, provided for use of newlib malloc.
+   Each module should have it's own heap allocated.
+   Beware that heap usage increases with OpenMP teams.  See also arenas.  */
+
+static size_t gcn_kernel_heap_size = DEFAULT_GCN_HEAP_SIZE;
+
+/* Flag to decide whether print to stderr information about what is going on.
+   Set in init_debug depending on environment variables.  */
+
+static bool debug;
+
+/* Flag to decide if the runtime should suppress a possible fallback to host
+   execution.  */
+
+static bool suppress_host_fallback;
+
+/* Flag to locate HSA runtime shared library that is dlopened
+   by this plug-in.  */
+
+static const char *hsa_runtime_lib;
+
+/* Flag to decide if the runtime should support also CPU devices (can be
+   a simulator).  */
+
+static bool support_cpu_devices;
+
+/* Runtime dimension overrides.  Zero indicates default.  */
+
+static int override_x_dim = 0;
+static int override_z_dim = 0;
+
+/* }}}  */
+/* {{{ Debug & Diagnostic  */
+
+/* Print a message to stderr if GCN_DEBUG value is set to true.  */
+
+#define DEBUG_PRINT(...) \
+  do \
+  { \
+    if (debug) \
+      { \
+ fprintf (stderr, __VA_ARGS__); \
+      } \
+  } \
+  while (false);
+
+/* Flush stderr if GCN_DEBUG value is set to true.  */
+
+#define DEBUG_FLUSH() \
+  do { \
+    if (debug) \
+      fflush (stderr); \
+  } while (false)
+
+/* Print a logging message with PREFIX to stderr if GCN_DEBUG value
+   is set to true.  */
+
+#define DEBUG_LOG(prefix, ...) \
+  do \
+    { \
+      DEBUG_PRINT (prefix); \
+      DEBUG_PRINT (__VA_ARGS__); \
+      DEBUG_FLUSH (); \
+    } while (false)
+
+/* Print a debugging message to stderr.  */
+
+#define GCN_DEBUG(...) DEBUG_LOG ("GCN debug: ", __VA_ARGS__)
+
+/* Print a warning message to stderr.  */
+
+#define GCN_WARNING(...) DEBUG_LOG ("GCN warning: ", __VA_ARGS__)
+
+/* Print HSA warning STR with an HSA STATUS code.  */
+
+static void
+hsa_warn (const char *str, hsa_status_t status)
+{
+  if (!debug)
+    return;
+
+  const char *hsa_error_msg = "[unknown]";
+  hsa_fns.hsa_status_string_fn (status, &hsa_error_msg);
+
+  fprintf (stderr, "GCN warning: %s\nRuntime message: %s\n", str,
+   hsa_error_msg);
+}
+
+/* Report a fatal error STR together with the HSA error corresponding to STATUS
+   and terminate execution of the current process.  */
+
+static void
+hsa_fatal (const char *str, hsa_status_t status)
+{
+  const char *hsa_error_msg = "[unknown]";
+  hsa_fns.hsa_status_string_fn (status, &hsa_error_msg);
+  GOMP_PLUGIN_fatal ("GCN fatal error: %s\nRuntime message: %s\n", str,
+     hsa_error_msg);
+}
+
+/* Like hsa_fatal, except only report error message, and return FALSE
+   for propagating error processing to outside of plugin.  */
+
+static bool
+hsa_error (const char *str, hsa_status_t status)
+{
+  const char *hsa_error_msg = "[unknown]";
+  hsa_fns.hsa_status_string_fn (status, &hsa_error_msg);
+  GOMP_PLUGIN_error ("GCN fatal error: %s\nRuntime message: %s\n", str,
+     hsa_error_msg);
+  return false;
+}
+
+/* Dump information about the available hardware.  */
+
+static void
+dump_hsa_system_info (void)
+{
+  hsa_status_t status;
+
+  hsa_endianness_t endianness;
+  status = hsa_fns.hsa_system_get_info_fn (HSA_SYSTEM_INFO_ENDIANNESS,
+   &endianness);
+  if (status == HSA_STATUS_SUCCESS)
+    switch (endianness)
+      {
+      case HSA_ENDIANNESS_LITTLE:
+ GCN_DEBUG ("HSA_SYSTEM_INFO_ENDIANNESS: LITTLE\n");
+ break;
+      case HSA_ENDIANNESS_BIG:
+ GCN_DEBUG ("HSA_SYSTEM_INFO_ENDIANNESS: BIG\n");
+ break;
+      default:
+ GCN_WARNING ("HSA_SYSTEM_INFO_ENDIANNESS: UNKNOWN\n");
+      }
+  else
+    GCN_WARNING ("HSA_SYSTEM_INFO_ENDIANNESS: FAILED\n");
+
+  uint8_t extensions[128];
+  status = hsa_fns.hsa_system_get_info_fn (HSA_SYSTEM_INFO_EXTENSIONS,
+   &extensions);
+  if (status == HSA_STATUS_SUCCESS)
+    {
+      if (extensions[0] & (1 << HSA_EXTENSION_IMAGES))
+ GCN_DEBUG ("HSA_SYSTEM_INFO_EXTENSIONS: IMAGES\n");
+    }
+  else
+    GCN_WARNING ("HSA_SYSTEM_INFO_EXTENSIONS: FAILED\n");
+}
+
+/* Dump information about the available hardware.  */
+
+static void
+dump_machine_model (hsa_machine_model_t machine_model, const char *s)
+{
+  switch (machine_model)
+    {
+    case HSA_MACHINE_MODEL_SMALL:
+      GCN_DEBUG ("%s: SMALL\n", s);
+      break;
+    case HSA_MACHINE_MODEL_LARGE:
+      GCN_DEBUG ("%s: LARGE\n", s);
+      break;
+    default:
+      GCN_WARNING ("%s: UNKNOWN\n", s);
+      break;
+    }
+}
+
+/* Dump information about the available hardware.  */
+
+static void
+dump_profile (hsa_profile_t profile, const char *s)
+{
+  switch (profile)
+    {
+    case HSA_PROFILE_FULL:
+      GCN_DEBUG ("%s: FULL\n", s);
+      break;
+    case HSA_PROFILE_BASE:
+      GCN_DEBUG ("%s: BASE\n", s);
+      break;
+    default:
+      GCN_WARNING ("%s: UNKNOWN\n", s);
+      break;
+    }
+}
+
+/* Dump information about a device memory region.  */
+
+static hsa_status_t
+dump_hsa_region (hsa_region_t region, void *data __attribute__((unused)))
+{
+  hsa_status_t status;
+
+  hsa_region_segment_t segment;
+  status = hsa_fns.hsa_region_get_info_fn (region, HSA_REGION_INFO_SEGMENT,
+   &segment);
+  if (status == HSA_STATUS_SUCCESS)
+    {
+      if (segment == HSA_REGION_SEGMENT_GLOBAL)
+ GCN_DEBUG ("HSA_REGION_INFO_SEGMENT: GLOBAL\n");
+      else if (segment == HSA_REGION_SEGMENT_READONLY)
+ GCN_DEBUG ("HSA_REGION_INFO_SEGMENT: READONLY\n");
+      else if (segment == HSA_REGION_SEGMENT_PRIVATE)
+ GCN_DEBUG ("HSA_REGION_INFO_SEGMENT: PRIVATE\n");
+      else if (segment == HSA_REGION_SEGMENT_GROUP)
+ GCN_DEBUG ("HSA_REGION_INFO_SEGMENT: GROUP\n");
+      else
+ GCN_WARNING ("HSA_REGION_INFO_SEGMENT: UNKNOWN\n");
+    }
+  else
+    GCN_WARNING ("HSA_REGION_INFO_SEGMENT: FAILED\n");
+
+  if (segment == HSA_REGION_SEGMENT_GLOBAL)
+    {
+      uint32_t flags;
+      status
+ = hsa_fns.hsa_region_get_info_fn (region, HSA_REGION_INFO_GLOBAL_FLAGS,
+  &flags);
+      if (status == HSA_STATUS_SUCCESS)
+ {
+  if (flags & HSA_REGION_GLOBAL_FLAG_KERNARG)
+    GCN_DEBUG ("HSA_REGION_INFO_GLOBAL_FLAGS: KERNARG\n");
+  if (flags & HSA_REGION_GLOBAL_FLAG_FINE_GRAINED)
+    GCN_DEBUG ("HSA_REGION_INFO_GLOBAL_FLAGS: FINE_GRAINED\n");
+  if (flags & HSA_REGION_GLOBAL_FLAG_COARSE_GRAINED)
+    GCN_DEBUG ("HSA_REGION_INFO_GLOBAL_FLAGS: COARSE_GRAINED\n");
+ }
+      else
+ GCN_WARNING ("HSA_REGION_INFO_GLOBAL_FLAGS: FAILED\n");
+    }
+
+  size_t size;
+  status = hsa_fns.hsa_region_get_info_fn (region, HSA_REGION_INFO_SIZE, &size);
+  if (status == HSA_STATUS_SUCCESS)
+    GCN_DEBUG ("HSA_REGION_INFO_SIZE: %zu\n", size);
+  else
+    GCN_WARNING ("HSA_REGION_INFO_SIZE: FAILED\n");
+
+  status
+    = hsa_fns.hsa_region_get_info_fn (region, HSA_REGION_INFO_ALLOC_MAX_SIZE,
+      &size);
+  if (status == HSA_STATUS_SUCCESS)
+    GCN_DEBUG ("HSA_REGION_INFO_ALLOC_MAX_SIZE: %zu\n", size);
+  else
+    GCN_WARNING ("HSA_REGION_INFO_ALLOC_MAX_SIZE: FAILED\n");
+
+  bool alloc_allowed;
+  status
+    = hsa_fns.hsa_region_get_info_fn (region,
+      HSA_REGION_INFO_RUNTIME_ALLOC_ALLOWED,
+      &alloc_allowed);
+  if (status == HSA_STATUS_SUCCESS)
+    GCN_DEBUG ("HSA_REGION_INFO_RUNTIME_ALLOC_ALLOWED: %u\n", alloc_allowed);
+  else
+    GCN_WARNING ("HSA_REGION_INFO_RUNTIME_ALLOC_ALLOWED: FAILED\n");
+
+  if (status != HSA_STATUS_SUCCESS || !alloc_allowed)
+    return HSA_STATUS_SUCCESS;
+
+  status
+    = hsa_fns.hsa_region_get_info_fn (region,
+      HSA_REGION_INFO_RUNTIME_ALLOC_GRANULE,
+      &size);
+  if (status == HSA_STATUS_SUCCESS)
+    GCN_DEBUG ("HSA_REGION_INFO_RUNTIME_ALLOC_GRANULE: %zu\n", size);
+  else
+    GCN_WARNING ("HSA_REGION_INFO_RUNTIME_ALLOC_GRANULE: FAILED\n");
+
+  size_t align;
+  status
+    = hsa_fns.hsa_region_get_info_fn (region,
+      HSA_REGION_INFO_RUNTIME_ALLOC_ALIGNMENT,
+      &align);
+  if (status == HSA_STATUS_SUCCESS)
+    GCN_DEBUG ("HSA_REGION_INFO_RUNTIME_ALLOC_ALIGNMENT: %zu\n", align);
+  else
+    GCN_WARNING ("HSA_REGION_INFO_RUNTIME_ALLOC_ALIGNMENT: FAILED\n");
+
+  return HSA_STATUS_SUCCESS;
+}
+
+/* Dump information about all the device memory regions.  */
+
+static void
+dump_hsa_regions (hsa_agent_t agent)
+{
+  hsa_status_t status;
+  status = hsa_fns.hsa_agent_iterate_regions_fn (agent,
+ dump_hsa_region,
+ NULL);
+  if (status != HSA_STATUS_SUCCESS)
+    hsa_error ("Dumping hsa regions failed", status);
+}
+
+/* Dump information about the available devices.  */
+
+static hsa_status_t
+dump_hsa_agent_info (hsa_agent_t agent, void *data __attribute__((unused)))
+{
+  hsa_status_t status;
+
+  char buf[64];
+  status = hsa_fns.hsa_agent_get_info_fn (agent, HSA_AGENT_INFO_NAME,
+  &buf);
+  if (status == HSA_STATUS_SUCCESS)
+    GCN_DEBUG ("HSA_AGENT_INFO_NAME: %s\n", buf);
+  else
+    GCN_WARNING ("HSA_AGENT_INFO_NAME: FAILED\n");
+
+  status = hsa_fns.hsa_agent_get_info_fn (agent, HSA_AGENT_INFO_VENDOR_NAME,
+  &buf);
+  if (status == HSA_STATUS_SUCCESS)
+    GCN_DEBUG ("HSA_AGENT_INFO_VENDOR_NAME: %s\n", buf);
+  else
+    GCN_WARNING ("HSA_AGENT_INFO_VENDOR_NAME: FAILED\n");
+
+  hsa_machine_model_t machine_model;
+  status
+    = hsa_fns.hsa_agent_get_info_fn (agent, HSA_AGENT_INFO_MACHINE_MODEL,
+     &machine_model);
+  if (status == HSA_STATUS_SUCCESS)
+    dump_machine_model (machine_model, "HSA_AGENT_INFO_MACHINE_MODEL");
+  else
+    GCN_WARNING ("HSA_AGENT_INFO_MACHINE_MODEL: FAILED\n");
+
+  hsa_profile_t profile;
+  status = hsa_fns.hsa_agent_get_info_fn (agent, HSA_AGENT_INFO_PROFILE,
+  &profile);
+  if (status == HSA_STATUS_SUCCESS)
+    dump_profile (profile, "HSA_AGENT_INFO_PROFILE");
+  else
+    GCN_WARNING ("HSA_AGENT_INFO_PROFILE: FAILED\n");
+
+  hsa_device_type_t device_type;
+  status = hsa_fns.hsa_agent_get_info_fn (agent, HSA_AGENT_INFO_DEVICE,
+  &device_type);
+  if (status == HSA_STATUS_SUCCESS)
+    {
+      switch (device_type)
+ {
+ case HSA_DEVICE_TYPE_CPU:
+  GCN_DEBUG ("HSA_AGENT_INFO_DEVICE: CPU\n");
+  break;
+ case HSA_DEVICE_TYPE_GPU:
+  GCN_DEBUG ("HSA_AGENT_INFO_DEVICE: GPU\n");
+  break;
+ case HSA_DEVICE_TYPE_DSP:
+  GCN_DEBUG ("HSA_AGENT_INFO_DEVICE: DSP\n");
+  break;
+ default:
+  GCN_WARNING ("HSA_AGENT_INFO_DEVICE: UNKNOWN\n");
+  break;
+ }
+    }
+  else
+    GCN_WARNING ("HSA_AGENT_INFO_DEVICE: FAILED\n");
+
+  uint32_t cu_count;
+  status = hsa_fns.hsa_agent_get_info_fn
+    (agent, HSA_AMD_AGENT_INFO_COMPUTE_UNIT_COUNT, &cu_count);
+  if (status == HSA_STATUS_SUCCESS)
+    GCN_DEBUG ("HSA_AMD_AGENT_INFO_COMPUTE_UNIT_COUNT: %u\n", cu_count);
+  else
+    GCN_WARNING ("HSA_AMD_AGENT_INFO_COMPUTE_UNIT_COUNT: FAILED\n");
+
+  uint32_t size;
+  status = hsa_fns.hsa_agent_get_info_fn (agent, HSA_AGENT_INFO_WAVEFRONT_SIZE,
+  &size);
+  if (status == HSA_STATUS_SUCCESS)
+    GCN_DEBUG ("HSA_AGENT_INFO_WAVEFRONT_SIZE: %u\n", size);
+  else
+    GCN_WARNING ("HSA_AGENT_INFO_WAVEFRONT_SIZE: FAILED\n");
+
+  uint32_t max_dim;
+  status = hsa_fns.hsa_agent_get_info_fn (agent,
+  HSA_AGENT_INFO_WORKGROUP_MAX_DIM,
+  &max_dim);
+  if (status == HSA_STATUS_SUCCESS)
+    GCN_DEBUG ("HSA_AGENT_INFO_WORKGROUP_MAX_DIM: %u\n", max_dim);
+  else
+    GCN_WARNING ("HSA_AGENT_INFO_WORKGROUP_MAX_DIM: FAILED\n");
+
+  uint32_t max_size;
+  status = hsa_fns.hsa_agent_get_info_fn (agent,
+  HSA_AGENT_INFO_WORKGROUP_MAX_SIZE,
+  &max_size);
+  if (status == HSA_STATUS_SUCCESS)
+    GCN_DEBUG ("HSA_AGENT_INFO_WORKGROUP_MAX_SIZE: %u\n", max_size);
+  else
+    GCN_WARNING ("HSA_AGENT_INFO_WORKGROUP_MAX_SIZE: FAILED\n");
+
+  uint32_t grid_max_dim;
+  status = hsa_fns.hsa_agent_get_info_fn (agent, HSA_AGENT_INFO_GRID_MAX_DIM,
+  &grid_max_dim);
+  if (status == HSA_STATUS_SUCCESS)
+    GCN_DEBUG ("HSA_AGENT_INFO_GRID_MAX_DIM: %u\n", grid_max_dim);
+  else
+    GCN_WARNING ("HSA_AGENT_INFO_GRID_MAX_DIM: FAILED\n");
+
+  uint32_t grid_max_size;
+  status = hsa_fns.hsa_agent_get_info_fn (agent, HSA_AGENT_INFO_GRID_MAX_SIZE,
+  &grid_max_size);
+  if (status == HSA_STATUS_SUCCESS)
+    GCN_DEBUG ("HSA_AGENT_INFO_GRID_MAX_SIZE: %u\n", grid_max_size);
+  else
+    GCN_WARNING ("HSA_AGENT_INFO_GRID_MAX_SIZE: FAILED\n");
+
+  dump_hsa_regions (agent);
+
+  return HSA_STATUS_SUCCESS;
+}
+
+/* Forward reference.  */
+
+static char *get_executable_symbol_name (hsa_executable_symbol_t symbol);
+
+/* Helper function for dump_executable_symbols.  */
+
+static hsa_status_t
+dump_executable_symbol (hsa_executable_t executable,
+ hsa_executable_symbol_t symbol,
+ void *data __attribute__((unused)))
+{
+  char *name = get_executable_symbol_name (symbol);
+
+  if (name)
+    {
+      GCN_DEBUG ("executable symbol: %s\n", name);
+      free (name);
+    }
+
+  return HSA_STATUS_SUCCESS;
+}
+
+/* Dump all global symbol in an executable.  */
+
+static void
+dump_executable_symbols (hsa_executable_t executable)
+{
+  hsa_status_t status;
+  status
+    = hsa_fns.hsa_executable_iterate_symbols_fn (executable,
+ dump_executable_symbol,
+ NULL);
+  if (status != HSA_STATUS_SUCCESS)
+    hsa_fatal ("Could not dump HSA executable symbols", status);
+}
+
+/* Dump kernel DISPATCH data structure and indent it by INDENT spaces.  */
+
+static void
+print_kernel_dispatch (struct kernel_dispatch *dispatch, unsigned indent)
+{
+  struct kernargs *kernargs = (struct kernargs *)dispatch->kernarg_address;
+
+  fprintf (stderr, "%*sthis: %p\n", indent, "", dispatch);
+  fprintf (stderr, "%*squeue: %p\n", indent, "", dispatch->queue);
+  fprintf (stderr, "%*skernarg_address: %p\n", indent, "", kernargs);
+  fprintf (stderr, "%*sheap address: %p\n", indent, "",
+   (void*)kernargs->heap_ptr);
+  fprintf (stderr, "%*sarena address: %p\n", indent, "",
+   (void*)kernargs->arena_ptr);
+  fprintf (stderr, "%*sobject: %lu\n", indent, "", dispatch->object);
+  fprintf (stderr, "%*sprivate_segment_size: %u\n", indent, "",
+   dispatch->private_segment_size);
+  fprintf (stderr, "%*sgroup_segment_size: %u\n", indent, "",
+   dispatch->group_segment_size);
+  fprintf (stderr, "\n");
+}
+
+/* }}}  */
+/* {{{ Utility functions  */
+
+/* Cast the thread local storage to gcn_thread.  */
+
+static inline struct gcn_thread *
+gcn_thread (void)
+{
+  return (struct gcn_thread *) GOMP_PLUGIN_acc_thread ();
+}
+
+/* Initialize debug and suppress_host_fallback according to the environment.  */
+
+static void
+init_environment_variables (void)
+{
+  if (secure_getenv ("GCN_DEBUG"))
+    debug = true;
+  else
+    debug = false;
+
+  if (secure_getenv ("GCN_SUPPRESS_HOST_FALLBACK"))
+    suppress_host_fallback = true;
+  else
+    suppress_host_fallback = false;
+
+  hsa_runtime_lib = secure_getenv ("HSA_RUNTIME_LIB");
+  if (hsa_runtime_lib == NULL)
+    hsa_runtime_lib = HSA_RUNTIME_LIB "libhsa-runtime64.so";
+
+  support_cpu_devices = secure_getenv ("GCN_SUPPORT_CPU_DEVICES");
+
+  const char *x = secure_getenv ("GCN_NUM_TEAMS");
+  if (!x)
+    x = secure_getenv ("GCN_NUM_GANGS");
+  if (x)
+    override_x_dim = atoi (x);
+
+  const char *z = secure_getenv ("GCN_NUM_THREADS");
+  if (!z)
+    z = secure_getenv ("GCN_NUM_WORKERS");
+  if (z)
+    override_z_dim = atoi (z);
+
+  const char *heap = secure_getenv ("GCN_HEAP_SIZE");
+  if (heap)
+    {
+      size_t tmp = atol (heap);
+      if (tmp)
+ gcn_kernel_heap_size = tmp;
+    }
+}
+
+/* Return malloc'd string with name of SYMBOL.  */
+
+static char *
+get_executable_symbol_name (hsa_executable_symbol_t symbol)
+{
+  hsa_status_t status;
+  char *res;
+  uint32_t len;
+  const hsa_executable_symbol_info_t info_name_length
+    = HSA_EXECUTABLE_SYMBOL_INFO_NAME_LENGTH;
+
+  status = hsa_fns.hsa_executable_symbol_get_info_fn (symbol, info_name_length,
+      &len);
+  if (status != HSA_STATUS_SUCCESS)
+    {
+      hsa_error ("Could not get length of symbol name", status);
+      return NULL;
+    }
+
+  res = GOMP_PLUGIN_malloc (len + 1);
+
+  const hsa_executable_symbol_info_t info_name
+    = HSA_EXECUTABLE_SYMBOL_INFO_NAME;
+
+  status = hsa_fns.hsa_executable_symbol_get_info_fn (symbol, info_name, res);
+
+  if (status != HSA_STATUS_SUCCESS)
+    {
+      hsa_error ("Could not get symbol name", status);
+      free (res);
+      return NULL;
+    }
+
+  res[len] = '\0';
+
+  return res;
+}
+
+/* Helper function for find_executable_symbol.  */
+
+static hsa_status_t
+find_executable_symbol_1 (hsa_executable_t executable,
+  hsa_executable_symbol_t symbol,
+  void *data)
+{
+  hsa_executable_symbol_t *res = (hsa_executable_symbol_t *)data;
+  *res = symbol;
+  return HSA_STATUS_INFO_BREAK;
+}
+
+/* Find a global symbol in EXECUTABLE, save to *SYMBOL and return true.  If not
+   found, return false.  */
+
+static bool
+find_executable_symbol (hsa_executable_t executable,
+ hsa_executable_symbol_t *symbol)
+{
+  hsa_status_t status;
+
+  status
+    = hsa_fns.hsa_executable_iterate_symbols_fn (executable,
+ find_executable_symbol_1,
+ symbol);
+  if (status != HSA_STATUS_INFO_BREAK)
+    {
+      hsa_error ("Could not find executable symbol", status);
+      return false;
+    }
+
+  return true;
+}
+
+/* Get the number of GPU Compute Units.  */
+
+static int
+get_cu_count (struct agent_info *agent)
+{
+  uint32_t cu_count;
+  hsa_status_t status = hsa_fns.hsa_agent_get_info_fn
+    (agent->id, HSA_AMD_AGENT_INFO_COMPUTE_UNIT_COUNT, &cu_count);
+  if (status == HSA_STATUS_SUCCESS)
+    return cu_count;
+  else
+    return 64;  /* The usual number for older devices.  */
+}
+
+/* Calculate the maximum grid size for OMP threads / OACC workers.
+   This depends on the kernel's resource usage levels.  */
+
+static int
+limit_worker_threads (int threads)
+{
+  /* FIXME Do something more inteligent here.
+     GCN can always run 4 threads within a Compute Unit, but
+     more than that depends on register usage.  */
+  if (threads > 16)
+    threads = 16;
+  return threads;
+}
+
+/* Parse the target attributes INPUT provided by the compiler and return true
+   if we should run anything all.  If INPUT is NULL, fill DEF with default
+   values, then store INPUT or DEF into *RESULT.
+
+   This is used for OpenMP only.  */
+
+static bool
+parse_target_attributes (void **input,
+ struct GOMP_kernel_launch_attributes *def,
+ struct GOMP_kernel_launch_attributes **result,
+ struct agent_info *agent)
+{
+  if (!input)
+    GOMP_PLUGIN_fatal ("No target arguments provided");
+
+  bool grid_attrs_found = false;
+  bool gcn_dims_found = false;
+  int gcn_teams = 0;
+  int gcn_threads = 0;
+  while (*input)
+    {
+      intptr_t id = (intptr_t) *input++, val;
+
+      if (id & GOMP_TARGET_ARG_SUBSEQUENT_PARAM)
+ val = (intptr_t) *input++;
+      else
+ val = id >> GOMP_TARGET_ARG_VALUE_SHIFT;
+
+      val = (val > INT_MAX) ? INT_MAX : val;
+
+      if ((id & GOMP_TARGET_ARG_DEVICE_MASK) == GOMP_DEVICE_GCN
+  && ((id & GOMP_TARGET_ARG_ID_MASK)
+      == GOMP_TARGET_ARG_HSA_KERNEL_ATTRIBUTES))
+ {
+  grid_attrs_found = true;
+  break;
+ }
+      else if ((id & GOMP_TARGET_ARG_DEVICE_ALL) == GOMP_TARGET_ARG_DEVICE_ALL)
+ {
+  gcn_dims_found = true;
+  switch (id & GOMP_TARGET_ARG_ID_MASK)
+    {
+    case GOMP_TARGET_ARG_NUM_TEAMS:
+      gcn_teams = val;
+      break;
+    case GOMP_TARGET_ARG_THREAD_LIMIT:
+      gcn_threads = limit_worker_threads (val);
+      break;
+    default:
+      ;
+    }
+ }
+    }
+
+  if (gcn_dims_found)
+    {
+      if (agent->gfx900_p && gcn_threads == 0 && override_z_dim == 0)
+ {
+  gcn_threads = 4;
+  GCN_WARNING ("VEGA BUG WORKAROUND: reducing default number of "
+       "threads to 4 per team.\n");
+  GCN_WARNING (" - If this is not a Vega 10 device, please use "
+       "GCN_NUM_THREADS=16\n");
+ }
+
+      def->ndim = 3;
+      /* Fiji has 64 CUs, but Vega20 has 60.  */
+      def->gdims[0] = (gcn_teams > 0) ? gcn_teams : get_cu_count (agent);
+      /* Each thread is 64 work items wide.  */
+      def->gdims[1] = 64;
+      /* A work group can have 16 wavefronts.  */
+      def->gdims[2] = (gcn_threads > 0) ? gcn_threads : 16;
+      def->wdims[0] = 1; /* Single team per work-group.  */
+      def->wdims[1] = 64;
+      def->wdims[2] = 16;
+      *result = def;
+      return true;
+    }
+  else if (!grid_attrs_found)
+    {
+      def->ndim = 1;
+      def->gdims[0] = 1;
+      def->gdims[1] = 1;
+      def->gdims[2] = 1;
+      def->wdims[0] = 1;
+      def->wdims[1] = 1;
+      def->wdims[2] = 1;
+      *result = def;
+      GCN_WARNING ("GOMP_OFFLOAD_run called with no launch attributes\n");
+      return true;
+    }
+
+  struct GOMP_kernel_launch_attributes *kla;
+  kla = (struct GOMP_kernel_launch_attributes *) *input;
+  *result = kla;
+  if (kla->ndim == 0 || kla->ndim > 3)
+    GOMP_PLUGIN_fatal ("Invalid number of dimensions (%u)", kla->ndim);
+
+  GCN_DEBUG ("GOMP_OFFLOAD_run called with %u dimensions:\n", kla->ndim);
+  unsigned i;
+  for (i = 0; i < kla->ndim; i++)
+    {
+      GCN_DEBUG ("  Dimension %u: grid size %u and group size %u\n", i,
+ kla->gdims[i], kla->wdims[i]);
+      if (kla->gdims[i] == 0)
+ return false;
+    }
+  return true;
+}
+
+/* Return the group size given the requested GROUP size, GRID size and number
+   of grid dimensions NDIM.  */
+
+static uint32_t
+get_group_size (uint32_t ndim, uint32_t grid, uint32_t group)
+{
+  if (group == 0)
+    {
+      /* TODO: Provide a default via environment or device characteristics.  */
+      if (ndim == 1)
+ group = 64;
+      else if (ndim == 2)
+ group = 8;
+      else
+ group = 4;
+    }
+
+  if (group > grid)
+    group = grid;
+  return group;
+}
+
+/* Atomically store pair of uint16_t values (HEADER and REST) to a PACKET.  */
+
+static void
+packet_store_release (uint32_t* packet, uint16_t header, uint16_t rest)
+{
+  __atomic_store_n (packet, header | (rest << 16), __ATOMIC_RELEASE);
+}
+
+/* A never-called callback for the HSA command queues.  These signal events
+   that we don't use, so we trigger an error.
+
+   This "queue" is not to be confused with the async queues, below.  */
+
+static void
+hsa_queue_callback (hsa_status_t status,
+ hsa_queue_t *queue __attribute__ ((unused)),
+ void *data __attribute__ ((unused)))
+{
+  hsa_fatal ("Asynchronous queue error", status);
+}
+
+/* }}}  */
+/* {{{ HSA initialization  */
+
+/* Populate hsa_fns with the function addresses from libhsa-runtime64.so.  */
+
+static bool
+init_hsa_runtime_functions (void)
+{
+#define DLSYM_FN(function) \
+  hsa_fns.function##_fn = dlsym (handle, #function); \
+  if (hsa_fns.function##_fn == NULL) \
+    return false;
+  void *handle = dlopen (hsa_runtime_lib, RTLD_LAZY);
+  if (handle == NULL)
+    return false;
+
+  DLSYM_FN (hsa_status_string)
+  DLSYM_FN (hsa_system_get_info)
+  DLSYM_FN (hsa_agent_get_info)
+  DLSYM_FN (hsa_init)
+  DLSYM_FN (hsa_iterate_agents)
+  DLSYM_FN (hsa_region_get_info)
+  DLSYM_FN (hsa_queue_create)
+  DLSYM_FN (hsa_agent_iterate_regions)
+  DLSYM_FN (hsa_executable_destroy)
+  DLSYM_FN (hsa_executable_create)
+  DLSYM_FN (hsa_executable_global_variable_define)
+  DLSYM_FN (hsa_executable_load_code_object)
+  DLSYM_FN (hsa_executable_freeze)
+  DLSYM_FN (hsa_signal_create)
+  DLSYM_FN (hsa_memory_allocate)
+  DLSYM_FN (hsa_memory_assign_agent)
+  DLSYM_FN (hsa_memory_copy)
+  DLSYM_FN (hsa_memory_free)
+  DLSYM_FN (hsa_signal_destroy)
+  DLSYM_FN (hsa_executable_get_symbol)
+  DLSYM_FN (hsa_executable_symbol_get_info)
+  DLSYM_FN (hsa_executable_iterate_symbols)
+  DLSYM_FN (hsa_queue_add_write_index_release)
+  DLSYM_FN (hsa_queue_load_read_index_acquire)
+  DLSYM_FN (hsa_signal_wait_acquire)
+  DLSYM_FN (hsa_signal_store_relaxed)
+  DLSYM_FN (hsa_signal_store_release)
+  DLSYM_FN (hsa_signal_load_acquire)
+  DLSYM_FN (hsa_queue_destroy)
+  DLSYM_FN (hsa_code_object_deserialize)
+  return true;
+#undef DLSYM_FN
+}
+
+/* Return true if the agent is a GPU and can accept of concurrent submissions
+   from different threads.  */
+
+static bool
+suitable_hsa_agent_p (hsa_agent_t agent)
+{
+  hsa_device_type_t device_type;
+  hsa_status_t status
+    = hsa_fns.hsa_agent_get_info_fn (agent, HSA_AGENT_INFO_DEVICE,
+     &device_type);
+  if (status != HSA_STATUS_SUCCESS)
+    return false;
+
+  switch (device_type)
+    {
+    case HSA_DEVICE_TYPE_GPU:
+      break;
+    case HSA_DEVICE_TYPE_CPU:
+      if (!support_cpu_devices)
+ return false;
+      break;
+    default:
+      return false;
+    }
+
+  uint32_t features = 0;
+  status = hsa_fns.hsa_agent_get_info_fn (agent, HSA_AGENT_INFO_FEATURE,
+  &features);
+  if (status != HSA_STATUS_SUCCESS
+      || !(features & HSA_AGENT_FEATURE_KERNEL_DISPATCH))
+    return false;
+  hsa_queue_type_t queue_type;
+  status = hsa_fns.hsa_agent_get_info_fn (agent, HSA_AGENT_INFO_QUEUE_TYPE,
+  &queue_type);
+  if (status != HSA_STATUS_SUCCESS
+      || (queue_type != HSA_QUEUE_TYPE_MULTI))
+    return false;
+
+  return true;
+}
+
+/* Callback of hsa_iterate_agents; if AGENT is a GPU device, increment
+   agent_count in hsa_context.  */
+
+static hsa_status_t
+count_gpu_agents (hsa_agent_t agent, void *data __attribute__ ((unused)))
+{
+  if (suitable_hsa_agent_p (agent))
+    hsa_context.agent_count++;
+  return HSA_STATUS_SUCCESS;
+}
+
+/* Callback of hsa_iterate_agents; if AGENT is a GPU device, assign the agent
+   id to the describing structure in the hsa context.  The index of the
+   structure is pointed to by DATA, increment it afterwards.  */
+
+static hsa_status_t
+assign_agent_ids (hsa_agent_t agent, void *data)
+{
+  if (suitable_hsa_agent_p (agent))
+    {
+      int *agent_index = (int *) data;
+      hsa_context.agents[*agent_index].id = agent;
+      ++*agent_index;
+    }
+  return HSA_STATUS_SUCCESS;
+}
+
+/* Initialize hsa_context if it has not already been done.
+   Return TRUE on success.  */
+
+static bool
+init_hsa_context (void)
+{
+  hsa_status_t status;
+  int agent_index = 0;
+
+  if (hsa_context.initialized)
+    return true;
+  init_environment_variables ();
+  if (!init_hsa_runtime_functions ())
+    {
+      GCN_WARNING ("Run-time could not be dynamically opened\n");
+      if (suppress_host_fallback)
+ GOMP_PLUGIN_fatal ("GCN host fallback has been suppressed");
+      return false;
+    }
+  status = hsa_fns.hsa_init_fn ();
+  if (status != HSA_STATUS_SUCCESS)
+    return hsa_error ("Run-time could not be initialized", status);
+  GCN_DEBUG ("HSA run-time initialized for GCN\n");
+
+  if (debug)
+    dump_hsa_system_info ();
+
+  status = hsa_fns.hsa_iterate_agents_fn (count_gpu_agents, NULL);
+  if (status != HSA_STATUS_SUCCESS)
+    return hsa_error ("GCN GPU devices could not be enumerated", status);
+  GCN_DEBUG ("There are %i GCN GPU devices.\n", hsa_context.agent_count);
+
+  hsa_context.agents
+    = GOMP_PLUGIN_malloc_cleared (hsa_context.agent_count
+  * sizeof (struct agent_info));
+  status = hsa_fns.hsa_iterate_agents_fn (assign_agent_ids, &agent_index);
+  if (agent_index != hsa_context.agent_count)
+    {
+      GOMP_PLUGIN_error ("Failed to assign IDs to all GCN agents");
+      return false;
+    }
+
+  if (debug)
+    {
+      status = hsa_fns.hsa_iterate_agents_fn (dump_hsa_agent_info, NULL);
+      if (status != HSA_STATUS_SUCCESS)
+ GOMP_PLUGIN_error ("Failed to list all HSA runtime agents");
+    }
+
+  hsa_context.initialized = true;
+  return true;
+}
+
+/* Verify that hsa_context has already been initialized and return the
+   agent_info structure describing device number N.  Return NULL on error.  */
+
+static struct agent_info *
+get_agent_info (int n)
+{
+  if (!hsa_context.initialized)
+    {
+      GOMP_PLUGIN_error ("Attempt to use uninitialized GCN context.");
+      return NULL;
+    }
+  if (n >= hsa_context.agent_count)
+    {
+      GOMP_PLUGIN_error ("Request to operate on non-existent GCN device %i", n);
+      return NULL;
+    }
+  if (!hsa_context.agents[n].initialized)
+    {
+      GOMP_PLUGIN_error ("Attempt to use an uninitialized GCN agent.");
+      return NULL;
+    }
+  return &hsa_context.agents[n];
+}
+
+/* Callback of hsa_agent_iterate_regions, via get_*_memory_region functions.
+
+   Selects (breaks at) a suitable region of type KIND.  */
+
+static hsa_status_t
+get_memory_region (hsa_region_t region, hsa_region_t *retval,
+   hsa_region_global_flag_t kind)
+{
+  hsa_status_t status;
+  hsa_region_segment_t segment;
+
+  status = hsa_fns.hsa_region_get_info_fn (region, HSA_REGION_INFO_SEGMENT,
+   &segment);
+  if (status != HSA_STATUS_SUCCESS)
+    return status;
+  if (segment != HSA_REGION_SEGMENT_GLOBAL)
+    return HSA_STATUS_SUCCESS;
+
+  uint32_t flags;
+  status = hsa_fns.hsa_region_get_info_fn (region, HSA_REGION_INFO_GLOBAL_FLAGS,
+   &flags);
+  if (status != HSA_STATUS_SUCCESS)
+    return status;
+  if (flags & kind)
+    {
+      *retval = region;
+      return HSA_STATUS_INFO_BREAK;
+    }
+  return HSA_STATUS_SUCCESS;
+}
+
+/* Callback of hsa_agent_iterate_regions.
+
+   Selects a kernargs memory region.  */
+
+static hsa_status_t
+get_kernarg_memory_region (hsa_region_t region, void *data)
+{
+  return get_memory_region (region, (hsa_region_t *)data,
+    HSA_REGION_GLOBAL_FLAG_KERNARG);
+}
+
+/* Callback of hsa_agent_iterate_regions.
+
+   Selects a coarse-grained memory region suitable for the heap and
+   offload data.  */
+
+static hsa_status_t
+get_data_memory_region (hsa_region_t region, void *data)
+{
+  return get_memory_region (region, (hsa_region_t *)data,
+    HSA_REGION_GLOBAL_FLAG_COARSE_GRAINED);
+}
+
+/* }}}  */
+/* {{{ Run  */
+
+/* Create or reuse a team arena.
+
+   Team arenas are used by OpenMP to avoid calling malloc multiple times
+   while setting up each team.  This is purely a performance optimization.
+
+   Allocating an arena also costs performance, albeit on the host side, so
+   this function will reuse an existing arena if a large enough one is idle.
+   The arena is released, but not deallocated, when the kernel exits.  */
+
+static void *
+get_team_arena (struct agent_info *agent, int num_teams)
+{
+  struct team_arena_list **next_ptr = &agent->team_arena_list;
+  struct team_arena_list *item;
+
+  for (item = *next_ptr; item; next_ptr = &item->next, item = item->next)
+    {
+      if (item->num_teams < num_teams)
+ continue;
+
+      if (pthread_mutex_trylock (&item->in_use))
+ continue;
+
+      return item->arena;
+    }
+
+  GCN_DEBUG ("Creating a new arena for %d teams\n", num_teams);
+
+  if (pthread_mutex_lock (&agent->team_arena_write_lock))
+    {
+      GOMP_PLUGIN_error ("Could not lock a GCN agent program mutex");
+      return false;
+    }
+  item = malloc (sizeof (*item));
+  item->num_teams = num_teams;
+  item->next = NULL;
+  *next_ptr = item;
+
+  if (pthread_mutex_init (&item->in_use, NULL))
+    {
+      GOMP_PLUGIN_error ("Failed to initialize a GCN team arena write mutex");
+      return false;
+    }
+  if (pthread_mutex_lock (&item->in_use))
+    {
+      GOMP_PLUGIN_error ("Could not lock a GCN agent program mutex");
+      return false;
+    }
+  if (pthread_mutex_unlock (&agent->team_arena_write_lock))
+    {
+      GOMP_PLUGIN_error ("Could not unlock a GCN agent program mutex");
+      return false;
+    }
+
+  const int TEAM_ARENA_SIZE = 64*1024;  /* Must match libgomp.h.  */
+  hsa_status_t status;
+  status = hsa_fns.hsa_memory_allocate_fn (agent->data_region,
+   TEAM_ARENA_SIZE*num_teams,
+   &item->arena);
+  if (status != HSA_STATUS_SUCCESS)
+    hsa_fatal ("Could not allocate memory for GCN kernel arena", status);
+  status = hsa_fns.hsa_memory_assign_agent_fn (item->arena, agent->id,
+       HSA_ACCESS_PERMISSION_RW);
+  if (status != HSA_STATUS_SUCCESS)
+    hsa_fatal ("Could not assign arena memory to device", status);
+
+  return item->arena;
+}
+
+/* Mark a team arena available for reuse.  */
+
+static void
+release_team_arena (struct agent_info* agent, void *arena)
+{
+  struct team_arena_list *item;
+
+  for (item = agent->team_arena_list; item; item = item->next)
+    {
+      if (item->arena == arena)
+ {
+  if (pthread_mutex_unlock (&item->in_use))
+    GOMP_PLUGIN_error ("Could not unlock a GCN agent program mutex");
+  return;
+ }
+    }
+  GOMP_PLUGIN_error ("Could not find a GCN arena to release.");
+}
+
+/* Clean up all the allocated team arenas.  */
+
+static bool
+destroy_team_arenas (struct agent_info *agent)
+{
+  struct team_arena_list *item, *next;
+
+  for (item = agent->team_arena_list; item; item = next)
+    {
+      next = item->next;
+      hsa_fns.hsa_memory_free_fn (item->arena);
+      if (pthread_mutex_destroy (&item->in_use))
+ {
+  GOMP_PLUGIN_error ("Failed to destroy a GCN team arena mutex");
+  return false;
+ }
+      free (item);
+    }
+  agent->team_arena_list = NULL;
+
+  return true;
+}
+
+/* Allocate memory on a specified device.  */
+
+static void *
+alloc_by_agent (struct agent_info *agent, size_t size)
+{
+  GCN_DEBUG ("Allocating %zu bytes on device %d\n", size, agent->device_id);
+
+  /* Zero-size allocations are invalid, so in order to return a valid pointer
+     we need to pass a valid size.  One source of zero-size allocations is
+     kernargs for kernels that have no inputs or outputs (the kernel may
+     only use console output, for example).  */
+  if (size == 0)
+    size = 4;
+
+  void *ptr;
+  hsa_status_t status = hsa_fns.hsa_memory_allocate_fn (agent->data_region,
+ size, &ptr);
+  if (status != HSA_STATUS_SUCCESS)
+    {
+      hsa_error ("Could not allocate device memory", status);
+      return NULL;
+    }
+
+  status = hsa_fns.hsa_memory_assign_agent_fn (ptr, agent->id,
+       HSA_ACCESS_PERMISSION_RW);
+  if (status != HSA_STATUS_SUCCESS)
+    {
+      hsa_error ("Could not assign data memory to device", status);
+      return NULL;
+    }
+
+  struct goacc_thread *thr = GOMP_PLUGIN_goacc_thread ();
+  bool profiling_dispatch_p
+    = __builtin_expect (thr != NULL && thr->prof_info != NULL, false);
+  if (profiling_dispatch_p)
+    {
+      acc_prof_info *prof_info = thr->prof_info;
+      acc_event_info data_event_info;
+      acc_api_info *api_info = thr->api_info;
+
+      prof_info->event_type = acc_ev_alloc;
+
+      data_event_info.data_event.event_type = prof_info->event_type;
+      data_event_info.data_event.valid_bytes
+ = _ACC_DATA_EVENT_INFO_VALID_BYTES;
+      data_event_info.data_event.parent_construct
+ = acc_construct_parallel;
+      data_event_info.data_event.implicit = 1;
+      data_event_info.data_event.tool_info = NULL;
+      data_event_info.data_event.var_name = NULL;
+      data_event_info.data_event.bytes = size;
+      data_event_info.data_event.host_ptr = NULL;
+      data_event_info.data_event.device_ptr = (void *) ptr;
+
+      api_info->device_api = acc_device_api_other;
+
+      GOMP_PLUGIN_goacc_profiling_dispatch (prof_info, &data_event_info,
+    api_info);
+    }
+
+  return ptr;
+}
+
+/* Create kernel dispatch data structure for given KERNEL, along with
+   the necessary device signals and memory allocations.  */
+
+static struct kernel_dispatch *
+create_kernel_dispatch (struct kernel_info *kernel, int num_teams)
+{
+  struct agent_info *agent = kernel->agent;
+  struct kernel_dispatch *shadow
+    = GOMP_PLUGIN_malloc_cleared (sizeof (struct kernel_dispatch));
+
+  shadow->agent = kernel->agent;
+  shadow->object = kernel->object;
+
+  hsa_signal_t sync_signal;
+  hsa_status_t status = hsa_fns.hsa_signal_create_fn (1, 0, NULL, &sync_signal);
+  if (status != HSA_STATUS_SUCCESS)
+    hsa_fatal ("Error creating the GCN sync signal", status);
+
+  shadow->signal = sync_signal.handle;
+  shadow->private_segment_size = kernel->private_segment_size;
+  shadow->group_segment_size = kernel->group_segment_size;
+
+  /* We expect kernels to request a single pointer, explicitly, and the
+     rest of struct kernargs, implicitly.  If they request anything else
+     then something is wrong.  */
+  if (kernel->kernarg_segment_size > 8)
+    {
+      GOMP_PLUGIN_fatal ("Unexpectedly large kernargs segment requested");
+      return NULL;
+    }
+
+  status = hsa_fns.hsa_memory_allocate_fn (agent->kernarg_region,
+   sizeof (struct kernargs),
+   &shadow->kernarg_address);
+  if (status != HSA_STATUS_SUCCESS)
+    hsa_fatal ("Could not allocate memory for GCN kernel arguments", status);
+  struct kernargs *kernargs = shadow->kernarg_address;
+
+  /* Zero-initialize the output_data (minimum needed).  */
+  kernargs->out_ptr = (int64_t)&kernargs->output_data;
+  kernargs->output_data.next_output = 0;
+  for (unsigned i = 0;
+       i < (sizeof (kernargs->output_data.queue)
+    / sizeof (kernargs->output_data.queue[0]));
+       i++)
+    kernargs->output_data.queue[i].written = 0;
+  kernargs->output_data.consumed = 0;
+
+  /* Pass in the heap location.  */
+  kernargs->heap_ptr = (int64_t)kernel->module->heap;
+
+  /* Create an arena.  */
+  if (kernel->kind == KIND_OPENMP)
+    kernargs->arena_ptr = (int64_t)get_team_arena (agent, num_teams);
+  else
+    kernargs->arena_ptr = 0;
+
+  /* Ensure we can recognize unset return values.  */
+  kernargs->output_data.return_value = 0xcafe0000;
+
+  return shadow;
+}
+
+/* Output any data written to console output from the kernel.  It is expected
+   that this function is polled during kernel execution.
+
+   We print all entries from the last item printed to the next entry without
+   a "written" flag.  If the "final" flag is set then it'll continue right to
+   the end.
+
+   The print buffer is circular, but the from and to locations don't wrap when
+   the buffer does, so the output limit is UINT_MAX.  The target blocks on
+   output when the buffer is full.  */
+
+static void
+console_output (struct kernel_info *kernel, struct kernargs *kernargs,
+ bool final)
+{
+  unsigned int limit = (sizeof (kernargs->output_data.queue)
+ / sizeof (kernargs->output_data.queue[0]));
+
+  unsigned int from = __atomic_load_n (&kernargs->output_data.consumed,
+       __ATOMIC_ACQUIRE);
+  unsigned int to = kernargs->output_data.next_output;
+
+  if (from > to)
+    {
+      /* Overflow.  */
+      if (final)
+ printf ("GCN print buffer overflowed.\n");
+      return;
+    }
+
+  unsigned int i;
+  for (i = from; i < to; i++)
+    {
+      struct printf_data *data = &kernargs->output_data.queue[i%limit];
+
+      if (!data->written && !final)
+ break;
+
+      switch (data->type)
+ {
+ case 0: printf ("%.128s%ld\n", data->msg, data->ivalue); break;
+ case 1: printf ("%.128s%f\n", data->msg, data->dvalue); break;
+ case 2: printf ("%.128s%.128s\n", data->msg, data->text); break;
+ case 3: printf ("%.128s%.128s", data->msg, data->text); break;
+ default: printf ("GCN print buffer error!\n"); break;
+ }
+      data->written = 0;
+      __atomic_store_n (&kernargs->output_data.consumed, i+1,
+ __ATOMIC_RELEASE);
+    }
+  fflush (stdout);
+}
+
+/* Release data structure created for a kernel dispatch in SHADOW argument,
+   and clean up the signal and memory allocations.  */
+
+static void
+release_kernel_dispatch (struct kernel_dispatch *shadow)
+{
+  GCN_DEBUG ("Released kernel dispatch: %p\n", shadow);
+
+  struct kernargs *kernargs = shadow->kernarg_address;
+  void *arena = (void *)kernargs->arena_ptr;
+  if (arena)
+    release_team_arena (shadow->agent, arena);
+
+  hsa_fns.hsa_memory_free_fn (shadow->kernarg_address);
+
+  hsa_signal_t s;
+  s.handle = shadow->signal;
+  hsa_fns.hsa_signal_destroy_fn (s);
+
+  free (shadow);
+}
+
+/* Extract the properties from a kernel binary.  */
+
+static void
+init_kernel_properties (struct kernel_info *kernel)
+{
+  hsa_status_t status;
+  struct agent_info *agent = kernel->agent;
+  hsa_executable_symbol_t kernel_symbol;
+  status = hsa_fns.hsa_executable_get_symbol_fn (agent->executable, NULL,
+ kernel->name, agent->id,
+ 0, &kernel_symbol);
+  if (status != HSA_STATUS_SUCCESS)
+    {
+      hsa_warn ("Could not find symbol for kernel in the code object", status);
+      fprintf (stderr, "not found name: '%s'\n", kernel->name);
+      dump_executable_symbols (agent->executable);
+      goto failure;
+    }
+  GCN_DEBUG ("Located kernel %s\n", kernel->name);
+  status = hsa_fns.hsa_executable_symbol_get_info_fn
+    (kernel_symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT, &kernel->object);
+  if (status != HSA_STATUS_SUCCESS)
+    hsa_fatal ("Could not extract a kernel object from its symbol", status);
+  status = hsa_fns.hsa_executable_symbol_get_info_fn
+    (kernel_symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_KERNARG_SEGMENT_SIZE,
+     &kernel->kernarg_segment_size);
+  if (status != HSA_STATUS_SUCCESS)
+    hsa_fatal ("Could not get info about kernel argument size", status);
+  status = hsa_fns.hsa_executable_symbol_get_info_fn
+    (kernel_symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_GROUP_SEGMENT_SIZE,
+     &kernel->group_segment_size);
+  if (status != HSA_STATUS_SUCCESS)
+    hsa_fatal ("Could not get info about kernel group segment size", status);
+  status = hsa_fns.hsa_executable_symbol_get_info_fn
+    (kernel_symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_PRIVATE_SEGMENT_SIZE,
+     &kernel->private_segment_size);
+  if (status != HSA_STATUS_SUCCESS)
+    hsa_fatal ("Could not get info about kernel private segment size",
+       status);
+
+  /* The kernel type is not known until something tries to launch it.  */
+  kernel->kind = KIND_UNKNOWN;
+
+  GCN_DEBUG ("Kernel structure for %s fully initialized with "
+     "following segment sizes: \n", kernel->name);
+  GCN_DEBUG ("  group_segment_size: %u\n",
+     (unsigned) kernel->group_segment_size);
+  GCN_DEBUG ("  private_segment_size: %u\n",
+     (unsigned) kernel->private_segment_size);
+  GCN_DEBUG ("  kernarg_segment_size: %u\n",
+     (unsigned) kernel->kernarg_segment_size);
+  return;
+
+failure:
+  kernel->initialization_failed = true;
+}
+
+/* Do all the work that is necessary before running KERNEL for the first time.
+   The function assumes the program has been created, finalized and frozen by
+   create_and_finalize_hsa_program.  */
+
+static void
+init_kernel (struct kernel_info *kernel)
+{
+  if (pthread_mutex_lock (&kernel->init_mutex))
+    GOMP_PLUGIN_fatal ("Could not lock a GCN kernel initialization mutex");
+  if (kernel->initialized)
+    {
+      if (pthread_mutex_unlock (&kernel->init_mutex))
+ GOMP_PLUGIN_fatal ("Could not unlock a GCN kernel initialization "
+   "mutex");
+
+      return;
+    }
+
+  init_kernel_properties (kernel);
+
+  if (!kernel->initialization_failed)
+    {
+      GCN_DEBUG ("\n");
+
+      kernel->initialized = true;
+    }
+  if (pthread_mutex_unlock (&kernel->init_mutex))
+    GOMP_PLUGIN_fatal ("Could not unlock a GCN kernel initialization "
+       "mutex");
+}
+
+/* Run KERNEL on its agent, pass VARS to it as arguments and take
+   launch attributes from KLA.
+  
+   MODULE_LOCKED indicates that the caller already holds the lock and
+   run_kernel need not lock it again.
+   If AQ is NULL then agent->sync_queue will be used.  */
+
+static void
+run_kernel (struct kernel_info *kernel, void *vars,
+    struct GOMP_kernel_launch_attributes *kla,
+    struct goacc_asyncqueue *aq, bool module_locked)
+{
+  GCN_DEBUG ("GCN launch on queue: %d:%d\n", kernel->agent->device_id,
+     (aq ? aq->id : 0));
+  GCN_DEBUG ("GCN launch attribs: gdims:[");
+  int i;
+  for (i = 0; i < kla->ndim; ++i)
+    {
+      if (i)
+ DEBUG_PRINT (", ");
+      DEBUG_PRINT ("%u", kla->gdims[i]);
+    }
+  DEBUG_PRINT ("], normalized gdims:[");
+  for (i = 0; i < kla->ndim; ++i)
+    {
+      if (i)
+ DEBUG_PRINT (", ");
+      DEBUG_PRINT ("%u", kla->gdims[i] / kla->wdims[i]);
+    }
+  DEBUG_PRINT ("], wdims:[");
+  for (i = 0; i < kla->ndim; ++i)
+    {
+      if (i)
+ DEBUG_PRINT (", ");
+      DEBUG_PRINT ("%u", kla->wdims[i]);
+    }
+  DEBUG_PRINT ("]\n");
+  DEBUG_FLUSH ();
+
+  struct agent_info *agent = kernel->agent;
+  if (!module_locked && pthread_rwlock_rdlock (&agent->module_rwlock))
+    GOMP_PLUGIN_fatal ("Unable to read-lock a GCN agent rwlock");
+
+  if (!agent->initialized)
+    GOMP_PLUGIN_fatal ("Agent must be initialized");
+
+  if (!kernel->initialized)
+    GOMP_PLUGIN_fatal ("Called kernel must be initialized");
+
+  hsa_queue_t *command_q = (aq ? aq->hsa_queue : kernel->agent->sync_queue);
+
+  uint64_t index
+    = hsa_fns.hsa_queue_add_write_index_release_fn (command_q, 1);
+  GCN_DEBUG ("Got AQL index %llu\n", (long long int) index);
+
+  /* Wait until the queue is not full before writing the packet.   */
+  while (index - hsa_fns.hsa_queue_load_read_index_acquire_fn (command_q)
+ >= command_q->size)
+    ;
+
+  /* Do not allow the dimensions to be overridden when running
+     constructors or destructors.  */
+  int override_x = kernel->kind == KIND_UNKNOWN ? 0 : override_x_dim;
+  int override_z = kernel->kind == KIND_UNKNOWN ? 0 : override_z_dim;
+
+  hsa_kernel_dispatch_packet_t *packet;
+  packet = ((hsa_kernel_dispatch_packet_t *) command_q->base_address)
+   + index % command_q->size;
+
+  memset (((uint8_t *) packet) + 4, 0, sizeof (*packet) - 4);
+  packet->grid_size_x = override_x ? : kla->gdims[0];
+  packet->workgroup_size_x = get_group_size (kla->ndim,
+     packet->grid_size_x,
+     kla->wdims[0]);
+
+  if (kla->ndim >= 2)
+    {
+      packet->grid_size_y = kla->gdims[1];
+      packet->workgroup_size_y = get_group_size (kla->ndim, kla->gdims[1],
+ kla->wdims[1]);
+    }
+  else
+    {
+      packet->grid_size_y = 1;
+      packet->workgroup_size_y = 1;
+    }
+
+  if (kla->ndim == 3)
+    {
+      packet->grid_size_z = limit_worker_threads (override_z
+  ? : kla->gdims[2]);
+      packet->workgroup_size_z = get_group_size (kla->ndim,
+ packet->grid_size_z,
+ kla->wdims[2]);
+    }
+  else
+    {
+      packet->grid_size_z = 1;
+      packet->workgroup_size_z = 1;
+    }
+
+  GCN_DEBUG ("GCN launch actuals: grid:[%u, %u, %u],"
+     " normalized grid:[%u, %u, %u], workgroup:[%u, %u, %u]\n",
+     packet->grid_size_x, packet->grid_size_y, packet->grid_size_z,
+     packet->grid_size_x / packet->workgroup_size_x,
+     packet->grid_size_y / packet->workgroup_size_y,
+     packet->grid_size_z / packet->workgroup_size_z,
+     packet->workgroup_size_x, packet->workgroup_size_y,
+     packet->workgroup_size_z);
+
+  struct kernel_dispatch *shadow
+    = create_kernel_dispatch (kernel, packet->grid_size_x);
+  shadow->queue = command_q;
+
+  if (debug)
+    {
+      fprintf (stderr, "\nKernel has following dependencies:\n");
+      print_kernel_dispatch (shadow, 2);
+    }
+
+  packet->private_segment_size = kernel->private_segment_size;
+  packet->group_segment_size = kernel->group_segment_size;
+  packet->kernel_object = kernel->object;
+  packet->kernarg_address = shadow->kernarg_address;
+  hsa_signal_t s;
+  s.handle = shadow->signal;
+  packet->completion_signal = s;
+  hsa_fns.hsa_signal_store_relaxed_fn (s, 1);
+  memcpy (shadow->kernarg_address, &vars, sizeof (vars));
+
+  GCN_DEBUG ("Copying kernel runtime pointer to kernarg_address\n");
+
+  uint16_t header;
+  header = HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE;
+  header |= HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE;
+  header |= HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE;
+
+  GCN_DEBUG ("Going to dispatch kernel %s on device %d\n", kernel->name,
+     agent->device_id);
+
+  packet_store_release ((uint32_t *) packet, header,
+ (uint16_t) kla->ndim
+ << HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS);
+
+  hsa_fns.hsa_signal_store_release_fn (command_q->doorbell_signal,
+       index);
+
+  GCN_DEBUG ("Kernel dispatched, waiting for completion\n");
+
+  /* Root signal waits with 1ms timeout.  */
+  while (hsa_fns.hsa_signal_wait_acquire_fn (s, HSA_SIGNAL_CONDITION_LT, 1,
+     1000 * 1000,
+     HSA_WAIT_STATE_BLOCKED) != 0)
+    {
+      console_output (kernel, shadow->kernarg_address, false);
+    }
+  console_output (kernel, shadow->kernarg_address, true);
+
+  struct kernargs *kernargs = shadow->kernarg_address;
+  unsigned int return_value = (unsigned int)kernargs->output_data.return_value;
+
+  release_kernel_dispatch (shadow);
+
+  if (!module_locked && pthread_rwlock_unlock (&agent->module_rwlock))
+    GOMP_PLUGIN_fatal ("Unable to unlock a GCN agent rwlock");
+
+  unsigned int upper = (return_value & ~0xffff) >> 16;
+  if (upper == 0xcafe)
+    ; // exit not called, normal termination.
+  else if (upper == 0xffff)
+    ; // exit called.
+  else
+    {
+      GOMP_PLUGIN_error ("Possible kernel exit value corruption, 2 most"
+ " significant bytes aren't 0xffff or 0xcafe: 0x%x\n",
+ return_value);
+      abort ();
+    }
+
+  if (upper == 0xffff)
+    {
+      unsigned int signal = (return_value >> 8) & 0xff;
+
+      if (signal == SIGABRT)
+ {
+  GCN_WARNING ("GCN Kernel aborted\n");
+  abort ();
+ }
+      else if (signal != 0)
+ {
+  GCN_WARNING ("GCN Kernel received unknown signal\n");
+  abort ();
+ }
+
+      GCN_DEBUG ("GCN Kernel exited with value: %d\n", return_value & 0xff);
+      exit (return_value & 0xff);
+    }
+}
+
+/* }}}  */
+/* {{{ Load/Unload  */
+
+/* Initialize KERNEL from D and other parameters.  Return true on success. */
+
+static bool
+init_basic_kernel_info (struct kernel_info *kernel,
+ struct hsa_kernel_description *d,
+ struct agent_info *agent,
+ struct module_info *module)
+{
+  kernel->agent = agent;
+  kernel->module = module;
+  kernel->name = d->name;
+  if (pthread_mutex_init (&kernel->init_mutex, NULL))
+    {
+      GOMP_PLUGIN_error ("Failed to initialize a GCN kernel mutex");
+      return false;
+    }
+  return true;
+}
+
+/* Find the load_offset for MODULE, save to *LOAD_OFFSET, and return true.  If
+   not found, return false.  */
+
+static bool
+find_load_offset (Elf64_Addr *load_offset, struct agent_info *agent,
+  struct module_info *module, Elf64_Ehdr *image,
+  Elf64_Shdr *sections)
+{
+  bool res = false;
+
+  hsa_status_t status;
+
+  hsa_executable_symbol_t symbol;
+  if (!find_executable_symbol (agent->executable, &symbol))
+    return false;
+
+  status = hsa_fns.hsa_executable_symbol_get_info_fn
+    (symbol, HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_ADDRESS, load_offset);
+  if (status != HSA_STATUS_SUCCESS)
+    {
+      hsa_error ("Could not extract symbol address", status);
+      return false;
+    }
+
+  char *symbol_name = get_executable_symbol_name (symbol);
+  if (symbol_name == NULL)
+    return false;
+
+  /* Find the kernel function in ELF, and calculate actual load offset.  */
+  for (int i = 0; i < image->e_shnum; i++)
+    if (sections[i].sh_type == SHT_SYMTAB)
+      {
+ Elf64_Shdr *strtab = &sections[sections[i].sh_link];
+ char *strings = (char *)image + strtab->sh_offset;
+
+ for (size_t offset = 0;
+     offset < sections[i].sh_size;
+     offset += sections[i].sh_entsize)
+  {
+    Elf64_Sym *sym = (Elf64_Sym*)((char*)image
+  + sections[i].sh_offset
+  + offset);
+    if (strcmp (symbol_name, strings + sym->st_name) == 0)
+      {
+ *load_offset -= sym->st_value;
+ res = true;
+ break;
+      }
+  }
+      }
+
+  free (symbol_name);
+  return res;
+}
+
+/* Create and finalize the program consisting of all loaded modules.  */
+
+static bool
+create_and_finalize_hsa_program (struct agent_info *agent)
+{
+  hsa_status_t status;
+  int reloc_count = 0;
+  bool res = true;
+  if (pthread_mutex_lock (&agent->prog_mutex))
+    {
+      GOMP_PLUGIN_error ("Could not lock a GCN agent program mutex");
+      return false;
+    }
+  if (agent->prog_finalized)
+    goto final;
+
+  status
+    = hsa_fns.hsa_executable_create_fn (HSA_PROFILE_FULL,
+ HSA_EXECUTABLE_STATE_UNFROZEN,
+ "", &agent->executable);
+  if (status != HSA_STATUS_SUCCESS)
+    {
+      hsa_error ("Could not create GCN executable", status);
+      goto fail;
+    }
+
+  /* Load any GCN modules.  */
+  struct module_info *module = agent->module;
+  if (module)
+    {
+      Elf64_Ehdr *image = (Elf64_Ehdr *)module->image_desc->gcn_image->image;
+
+      /* Hide relocations from the HSA runtime loader.
+ Keep a copy of the unmodified section headers to use later.  */
+      Elf64_Shdr *image_sections = (Elf64_Shdr *)((char *)image
+  + image->e_shoff);
+      for (int i = image->e_shnum - 1; i >= 0; i--)
+ {
+  if (image_sections[i].sh_type == SHT_RELA
+      || image_sections[i].sh_type == SHT_REL)
+    /* Change section type to something harmless.  */
+    image_sections[i].sh_type |= 0x80;
+ }
+
+      hsa_code_object_t co = { 0 };
+      status = hsa_fns.hsa_code_object_deserialize_fn
+ (module->image_desc->gcn_image->image,
+ module->image_desc->gcn_image->size,
+ NULL, &co);
+      if (status != HSA_STATUS_SUCCESS)
+ {
+  hsa_error ("Could not deserialize GCN code object", status);
+  goto fail;
+ }
+
+      status = hsa_fns.hsa_executable_load_code_object_fn
+ (agent->executable, agent->id, co, "");
+      if (status != HSA_STATUS_SUCCESS)
+ {
+  hsa_error ("Could not load GCN code object", status);
+  goto fail;
+ }
+
+      if (!module->heap)
+ {
+  status = hsa_fns.hsa_memory_allocate_fn (agent->data_region,
+   gcn_kernel_heap_size,
+   (void**)&module->heap);
+  if (status != HSA_STATUS_SUCCESS)
+    {
+      hsa_error ("Could not allocate memory for GCN heap", status);
+      goto fail;
+    }
+
+  status = hsa_fns.hsa_memory_assign_agent_fn
+ (module->heap, agent->id, HSA_ACCESS_PERMISSION_RW);
+  if (status != HSA_STATUS_SUCCESS)
+    {
+      hsa_error ("Could not assign GCN heap memory to device", status);
+      goto fail;
+    }
+
+  hsa_fns.hsa_memory_copy_fn (&module->heap->size,
+      &gcn_kernel_heap_size,
+      sizeof (gcn_kernel_heap_size));
+ }
+
+    }
+
+  if (debug)
+    dump_executable_symbols (agent->executable);
+
+  status = hsa_fns.hsa_executable_freeze_fn (agent->executable, "");
+  if (status != HSA_STATUS_SUCCESS)
+    {
+      hsa_error ("Could not freeze the GCN executable", status);
+      goto fail;
+    }
+
+  if (agent->module)
+    {
+      struct module_info *module = agent->module;
+      Elf64_Ehdr *image = (Elf64_Ehdr *)module->image_desc->gcn_image->image;
+      Elf64_Shdr *sections = (Elf64_Shdr *)((char *)image + image->e_shoff);
+
+      Elf64_Addr load_offset;
+      if (!find_load_offset (&load_offset, agent, module, image, sections))
+ goto fail;
+
+      /* Record the physical load address range.
+ We need this for data copies later.  */
+      Elf64_Phdr *segments = (Elf64_Phdr *)((char*)image + image->e_phoff);
+      Elf64_Addr low = ~0, high = 0;
+      for (int i = 0; i < image->e_phnum; i++)
+ if (segments[i].p_memsz > 0)
+  {
+    if (segments[i].p_paddr < low)
+      low = segments[i].p_paddr;
+    if (segments[i].p_paddr > high)
+      high = segments[i].p_paddr + segments[i].p_memsz - 1;
+  }
+      module->phys_address_start = low + load_offset;
+      module->phys_address_end = high + load_offset;
+
+      // Find dynamic symbol table
+      Elf64_Shdr *dynsym = NULL;
+      for (int i = 0; i < image->e_shnum; i++)
+ if (sections[i].sh_type == SHT_DYNSYM)
+  {
+    dynsym = &sections[i];
+    break;
+  }
+
+      /* Fix up relocations.  */
+      for (int i = 0; i < image->e_shnum; i++)
+ {
+  if (sections[i].sh_type == (SHT_RELA | 0x80))
+    for (size_t offset = 0;
+ offset < sections[i].sh_size;
+ offset += sections[i].sh_entsize)
+      {
+ Elf64_Rela *reloc = (Elf64_Rela*)((char*)image
+  + sections[i].sh_offset
+  + offset);
+ Elf64_Sym *sym =
+  (dynsym
+   ? (Elf64_Sym*)((char*)image
+  + dynsym->sh_offset
+  + (dynsym->sh_entsize
+     * ELF64_R_SYM (reloc->r_info)))
+   : NULL);
+
+ int64_t S = (sym ? sym->st_value : 0);
+ int64_t P = reloc->r_offset + load_offset;
+ int64_t A = reloc->r_addend;
+ int64_t B = load_offset;
+ int64_t V, size;
+ switch (ELF64_R_TYPE (reloc->r_info))
+  {
+  case R_AMDGPU_ABS32_LO:
+    V = (S + A) & 0xFFFFFFFF;
+    size = 4;
+    break;
+  case R_AMDGPU_ABS32_HI:
+    V = (S + A) >> 32;
+    size = 4;
+    break;
+  case R_AMDGPU_ABS64:
+    V = S + A;
+    size = 8;
+    break;
+  case R_AMDGPU_REL32:
+    V = S + A - P;
+    size = 4;
+    break;
+  case R_AMDGPU_REL64:
+    /* FIXME
+       LLD seems to emit REL64 where the the assembler has
+       ABS64.  This is clearly wrong because it's not what the
+       compiler is expecting.  Let's assume, for now, that
+       it's a bug.  In any case, GCN kernels are always self
+       contained and therefore relative relocations will have
+       been resolved already, so this should be a safe
+       workaround.  */
+    V = S + A/* - P*/;
+    size = 8;
+    break;
+  case R_AMDGPU_ABS32:
+    V = S + A;
+    size = 4;
+    break;
+    /* TODO R_AMDGPU_GOTPCREL */
+    /* TODO R_AMDGPU_GOTPCREL32_LO */
+    /* TODO R_AMDGPU_GOTPCREL32_HI */
+  case R_AMDGPU_REL32_LO:
+    V = (S + A - P) & 0xFFFFFFFF;
+    size = 4;
+    break;
+  case R_AMDGPU_REL32_HI:
+    V = (S + A - P) >> 32;
+    size = 4;
+    break;
+  case R_AMDGPU_RELATIVE64:
+    V = B + A;
+    size = 8;
+    break;
+  default:
+    fprintf (stderr, "Error: unsupported relocation type.\n");
+    exit (1);
+  }
+ status = hsa_fns.hsa_memory_copy_fn ((void*)P, &V, size);
+ if (status != HSA_STATUS_SUCCESS)
+  {
+    hsa_error ("Failed to fix up relocation", status);
+    goto fail;
+  }
+ reloc_count++;
+      }
+ }
+    }
+
+  GCN_DEBUG ("Loaded GCN kernels to device %d (%d relocations)\n",
+     agent->device_id, reloc_count);
+
+final:
+  agent->prog_finalized = true;
+
+  if (pthread_mutex_unlock (&agent->prog_mutex))
+    {
+      GOMP_PLUGIN_error ("Could not unlock a GCN agent program mutex");
+      res = false;
+    }
+
+  return res;
+
+fail:
+  res = false;
+  goto final;
+}
+
+/* Free the HSA program in agent and everything associated with it and set
+   agent->prog_finalized and the initialized flags of all kernels to false.
+   Return TRUE on success.  */
+
+static bool
+destroy_hsa_program (struct agent_info *agent)
+{
+  if (!agent->prog_finalized)
+    return true;
+
+  hsa_status_t status;
+
+  GCN_DEBUG ("Destroying the current GCN program.\n");
+
+  status = hsa_fns.hsa_executable_destroy_fn (agent->executable);
+  if (status != HSA_STATUS_SUCCESS)
+    return hsa_error ("Could not destroy GCN executable", status);
+
+  if (agent->module)
+    {
+      int i;
+      for (i = 0; i < agent->module->kernel_count; i++)
+ agent->module->kernels[i].initialized = false;
+
+      if (agent->module->heap)
+ {
+  hsa_fns.hsa_memory_free_fn (agent->module->heap);
+  agent->module->heap = NULL;
+ }
+    }
+  agent->prog_finalized = false;
+  return true;
+}
+
+/* Deinitialize all information associated with MODULE and kernels within
+   it.  Return TRUE on success.  */
+
+static bool
+destroy_module (struct module_info *module, bool locked)
+{
+  /* Run destructors before destroying module.  */
+  struct GOMP_kernel_launch_attributes kla =
+    { 3,
+      /* Grid size.  */
+      { 1, 64, 1 },
+      /* Work-group size.  */
+      { 1, 64, 1 }
+    };
+
+  if (module->fini_array_func)
+    {
+      init_kernel (module->fini_array_func);
+      run_kernel (module->fini_array_func, NULL, &kla, NULL, locked);
+    }
+  module->constructors_run_p = false;
+
+  int i;
+  for (i = 0; i < module->kernel_count; i++)
+    if (pthread_mutex_destroy (&module->kernels[i].init_mutex))
+      {
+ GOMP_PLUGIN_error ("Failed to destroy a GCN kernel initialization "
+   "mutex");
+ return false;
+      }
+
+  return true;
+}
+
+/* }}}  */
+/* {{{ Async  */
+
+/* Callback of dispatch queues to report errors.  */
+
+static void
+execute_queue_entry (struct goacc_asyncqueue *aq, int index)
+{
+  struct queue_entry *entry = &aq->queue[index];
+
+  switch (entry->type)
+    {
+    case KERNEL_LAUNCH:
+      if (DEBUG_QUEUES)
+ GCN_DEBUG ("Async thread %d:%d: Executing launch entry (%d)\n",
+   aq->agent->device_id, aq->id, index);
+      run_kernel (entry->u.launch.kernel,
+  entry->u.launch.vars,
+  &entry->u.launch.kla, aq, false);
+      if (DEBUG_QUEUES)
+ GCN_DEBUG ("Async thread %d:%d: Executing launch entry (%d) done\n",
+   aq->agent->device_id, aq->id, index);
+      break;
+
+    case CALLBACK:
+      if (DEBUG_QUEUES)
+ GCN_DEBUG ("Async thread %d:%d: Executing callback entry (%d)\n",
+   aq->agent->device_id, aq->id, index);
+      entry->u.callback.fn (entry->u.callback.data);
+      if (DEBUG_QUEUES)
+ GCN_DEBUG ("Async thread %d:%d: Executing callback entry (%d) done\n",
+   aq->agent->device_id, aq->id, index);
+      break;
+
+    case ASYNC_WAIT:
+      {
+ /* FIXME: is it safe to access a placeholder that may already have
+   been executed?  */
+        struct placeholder *placeholderp = entry->u.asyncwait.placeholderp;
+
+ if (DEBUG_QUEUES)
+          GCN_DEBUG ("Async thread %d:%d: Executing async wait entry (%d)\n",
+     aq->agent->device_id, aq->id, index);
+
+ pthread_mutex_lock (&placeholderp->mutex);
+
+ while (!placeholderp->executed)
+          pthread_cond_wait (&placeholderp->cond, &placeholderp->mutex);
+
+ pthread_mutex_unlock (&placeholderp->mutex);
+
+ if (pthread_cond_destroy (&placeholderp->cond))
+  GOMP_PLUGIN_error ("Failed to destroy serialization cond");
+
+ if (pthread_mutex_destroy (&placeholderp->mutex))
+  GOMP_PLUGIN_error ("Failed to destroy serialization mutex");
+
+ if (DEBUG_QUEUES)
+          GCN_DEBUG ("Async thread %d:%d: Executing async wait "
+     "entry (%d) done\n", aq->agent->device_id, aq->id, index);
+      }
+      break;
+
+    case ASYNC_PLACEHOLDER:
+      pthread_mutex_lock (&entry->u.placeholder.mutex);
+      entry->u.placeholder.executed = 1;
+      pthread_cond_signal (&entry->u.placeholder.cond);
+      pthread_mutex_unlock (&entry->u.placeholder.mutex);
+      break;
+
+    default:
+      GOMP_PLUGIN_fatal ("Unknown queue element");
+    }
+}
+
+/* This function is run as a thread to service an async queue in the
+   background.  It runs continuously until the stop flag is set.  */
+
+static void *
+drain_queue (void *thread_arg)
+{
+  struct goacc_asyncqueue *aq = thread_arg;
+
+  if (DRAIN_QUEUE_SYNCHRONOUS_P)
+    {
+      aq->drain_queue_stop = 2;
+      return NULL;
+    }
+
+  pthread_mutex_lock (&aq->mutex);
+
+  while (true)
+    {
+      if (aq->drain_queue_stop)
+ break;
+
+      if (aq->queue_n > 0)
+ {
+  pthread_mutex_unlock (&aq->mutex);
+  execute_queue_entry (aq, aq->queue_first);
+
+  pthread_mutex_lock (&aq->mutex);
+  aq->queue_first = ((aq->queue_first + 1)
+     % ASYNC_QUEUE_SIZE);
+  aq->queue_n--;
+
+  if (DEBUG_THREAD_SIGNAL)
+    GCN_DEBUG ("Async thread %d:%d: broadcasting queue out update\n",
+       aq->agent->device_id, aq->id);
+  pthread_cond_broadcast (&aq->queue_cond_out);
+  pthread_mutex_unlock (&aq->mutex);
+
+  if (DEBUG_QUEUES)
+    GCN_DEBUG ("Async thread %d:%d: continue\n", aq->agent->device_id,
+       aq->id);
+  pthread_mutex_lock (&aq->mutex);
+ }
+      else
+ {
+  if (DEBUG_THREAD_SLEEP)
+    GCN_DEBUG ("Async thread %d:%d: going to sleep\n",
+       aq->agent->device_id, aq->id);
+  pthread_cond_wait (&aq->queue_cond_in, &aq->mutex);
+  if (DEBUG_THREAD_SLEEP)
+    GCN_DEBUG ("Async thread %d:%d: woke up, rechecking\n",
+       aq->agent->device_id, aq->id);
+ }
+    }
+
+  aq->drain_queue_stop = 2;
+  if (DEBUG_THREAD_SIGNAL)
+    GCN_DEBUG ("Async thread %d:%d: broadcasting last queue out update\n",
+       aq->agent->device_id, aq->id);
+  pthread_cond_broadcast (&aq->queue_cond_out);
+  pthread_mutex_unlock (&aq->mutex);
+
+  GCN_DEBUG ("Async thread %d:%d: returning\n", aq->agent->device_id, aq->id);
+  return NULL;
+}
+
+/* This function is used only when DRAIN_QUEUE_SYNCHRONOUS_P is set, which
+   is not usually the case.  This is just a debug tool.  */
+
+static void
+drain_queue_synchronous (struct goacc_asyncqueue *aq)
+{
+  pthread_mutex_lock (&aq->mutex);
+
+  while (aq->queue_n > 0)
+    {
+      execute_queue_entry (aq, aq->queue_first);
+
+      aq->queue_first = ((aq->queue_first + 1)
+ % ASYNC_QUEUE_SIZE);
+      aq->queue_n--;
+    }
+
+  pthread_mutex_unlock (&aq->mutex);
+}
+
+/* Block the current thread until an async queue is writable.  */
+
+static void
+wait_for_queue_nonfull (struct goacc_asyncqueue *aq)
+{
+  if (aq->queue_n == ASYNC_QUEUE_SIZE)
+    {
+      pthread_mutex_lock (&aq->mutex);
+
+      /* Queue is full.  Wait for it to not be full.  */
+      while (aq->queue_n == ASYNC_QUEUE_SIZE)
+ pthread_cond_wait (&aq->queue_cond_out, &aq->mutex);
+
+      pthread_mutex_unlock (&aq->mutex);
+    }
+}
+
+/* Request an asynchronous kernel launch on the specified queue.  This
+   may block if the queue is full, but returns without waiting for the
+   kernel to run.  */
+
+static void
+queue_push_launch (struct goacc_asyncqueue *aq, struct kernel_info *kernel,
+   void *vars, struct GOMP_kernel_launch_attributes *kla)
+{
+  assert (aq->agent == kernel->agent);
+
+  wait_for_queue_nonfull (aq);
+
+  pthread_mutex_lock (&aq->mutex);
+
+  int queue_last = ((aq->queue_first + aq->queue_n)
+    % ASYNC_QUEUE_SIZE);
+  if (DEBUG_QUEUES)
+    GCN_DEBUG ("queue_push_launch %d:%d: at %i\n", aq->agent->device_id,
+       aq->id, queue_last);
+
+  aq->queue[queue_last].type = KERNEL_LAUNCH;
+  aq->queue[queue_last].u.launch.kernel = kernel;
+  aq->queue[queue_last].u.launch.vars = vars;
+  aq->queue[queue_last].u.launch.kla = *kla;
+
+  aq->queue_n++;
+
+  if (DEBUG_THREAD_SIGNAL)
+    GCN_DEBUG ("signalling async thread %d:%d: cond_in\n",
+       aq->agent->device_id, aq->id);
+  pthread_cond_signal (&aq->queue_cond_in);
+
+  pthread_mutex_unlock (&aq->mutex);
+}
+
+/* Request an asynchronous callback on the specified queue.  The callback
+   function will be called, with the given opaque data, from the appropriate
+   async thread, when all previous items on that queue are complete.  */
+
+static void
+queue_push_callback (struct goacc_asyncqueue *aq, void (*fn)(void *),
+     void *data)
+{
+  wait_for_queue_nonfull (aq);
+
+  pthread_mutex_lock (&aq->mutex);
+
+  int queue_last = ((aq->queue_first + aq->queue_n)
+    % ASYNC_QUEUE_SIZE);
+  if (DEBUG_QUEUES)
+    GCN_DEBUG ("queue_push_callback %d:%d: at %i\n", aq->agent->device_id,
+       aq->id, queue_last);
+
+  aq->queue[queue_last].type = CALLBACK;
+  aq->queue[queue_last].u.callback.fn = fn;
+  aq->queue[queue_last].u.callback.data = data;
+
+  aq->queue_n++;
+
+  if (DEBUG_THREAD_SIGNAL)
+    GCN_DEBUG ("signalling async thread %d:%d: cond_in\n",
+       aq->agent->device_id, aq->id);
+  pthread_cond_signal (&aq->queue_cond_in);
+
+  pthread_mutex_unlock (&aq->mutex);
+}
+
+/* Request that a given async thread wait for another thread (unspecified) to
+   reach the given placeholder.  The wait will occur when all previous entries
+   on the queue are complete.  A placeholder is effectively a kind of signal
+   which simply sets a flag when encountered in a queue.  */
+
+static void
+queue_push_asyncwait (struct goacc_asyncqueue *aq,
+      struct placeholder *placeholderp)
+{
+  wait_for_queue_nonfull (aq);
+
+  pthread_mutex_lock (&aq->mutex);
+
+  int queue_last = ((aq->queue_first + aq->queue_n) % ASYNC_QUEUE_SIZE);
+  if (DEBUG_QUEUES)
+    GCN_DEBUG ("queue_push_asyncwait %d:%d: at %i\n", aq->agent->device_id,
+       aq->id, queue_last);
+
+  aq->queue[queue_last].type = ASYNC_WAIT;
+  aq->queue[queue_last].u.asyncwait.placeholderp = placeholderp;
+
+  aq->queue_n++;
+
+  if (DEBUG_THREAD_SIGNAL)
+    GCN_DEBUG ("signalling async thread %d:%d: cond_in\n",
+       aq->agent->device_id, aq->id);
+  pthread_cond_signal (&aq->queue_cond_in);
+
+  pthread_mutex_unlock (&aq->mutex);
+}
+
+/* Add a placeholder into an async queue.  When the async thread reaches the
+   placeholder it will set the "executed" flag to true and continue.
+   Another thread may be waiting on this thread reaching the placeholder.  */
+
+static struct placeholder *
+queue_push_placeholder (struct goacc_asyncqueue *aq)
+{
+  struct placeholder *placeholderp;
+
+  wait_for_queue_nonfull (aq);
+
+  pthread_mutex_lock (&aq->mutex);
+
+  int queue_last = ((aq->queue_first + aq->queue_n) % ASYNC_QUEUE_SIZE);
+  if (DEBUG_QUEUES)
+    GCN_DEBUG ("queue_push_placeholder %d:%d: at %i\n", aq->agent->device_id,
+       aq->id, queue_last);
+
+  aq->queue[queue_last].type = ASYNC_PLACEHOLDER;
+  placeholderp = &aq->queue[queue_last].u.placeholder;
+
+  if (pthread_mutex_init (&placeholderp->mutex, NULL))
+    {
+      pthread_mutex_unlock (&aq->mutex);
+      GOMP_PLUGIN_error ("Failed to initialize serialization mutex");
+    }
+
+  if (pthread_cond_init (&placeholderp->cond, NULL))
+    {
+      pthread_mutex_unlock (&aq->mutex);
+      GOMP_PLUGIN_error ("Failed to initialize serialization cond");
+    }
+
+  placeholderp->executed = 0;
+
+  aq->queue_n++;
+
+  if (DEBUG_THREAD_SIGNAL)
+    GCN_DEBUG ("signalling async thread %d:%d: cond_in\n",
+       aq->agent->device_id, aq->id);
+  pthread_cond_signal (&aq->queue_cond_in);
+
+  pthread_mutex_unlock (&aq->mutex);
+
+  return placeholderp;
+}
+
+/* Signal an asynchronous thread to terminate, and wait for it to do so.  */
+
+static void
+finalize_async_thread (struct goacc_asyncqueue *aq)
+{
+  pthread_mutex_lock (&aq->mutex);
+  if (aq->drain_queue_stop == 2)
+    {
+      pthread_mutex_unlock (&aq->mutex);
+      return;
+    }
+
+  aq->drain_queue_stop = 1;
+
+  if (DEBUG_THREAD_SIGNAL)
+    GCN_DEBUG ("Signalling async thread %d:%d: cond_in\n",
+       aq->agent->device_id, aq->id);
+  pthread_cond_signal (&aq->queue_cond_in);
+
+  while (aq->drain_queue_stop != 2)
+    {
+      if (DEBUG_THREAD_SLEEP)
+ GCN_DEBUG ("Waiting for async thread %d:%d to finish, putting thread"
+   " to sleep\n", aq->agent->device_id, aq->id);
+      pthread_cond_wait (&aq->queue_cond_out, &aq->mutex);
+      if (DEBUG_THREAD_SLEEP)
+ GCN_DEBUG ("Waiting, woke up thread %d:%d.  Rechecking\n",
+   aq->agent->device_id, aq->id);
+    }
+
+  GCN_DEBUG ("Done waiting for async thread %d:%d\n", aq->agent->device_id,
+     aq->id);
+  pthread_mutex_unlock (&aq->mutex);
+
+  int err = pthread_join (aq->thread_drain_queue, NULL);
+  if (err != 0)
+    GOMP_PLUGIN_fatal ("Join async thread %d:%d: failed: %s",
+       aq->agent->device_id, aq->id, strerror (err));
+  GCN_DEBUG ("Joined with async thread %d:%d\n", aq->agent->device_id, aq->id);
+}
+
+/* Set up an async queue for OpenMP.  There will be only one.  The
+   implementation simply uses an OpenACC async queue.
+   FIXME: is this thread-safe if two threads call this function?  */
+
+static void
+maybe_init_omp_async (struct agent_info *agent)
+{
+  if (!agent->omp_async_queue)
+    agent->omp_async_queue
+      = GOMP_OFFLOAD_openacc_async_construct (agent->device_id);
+}
+
+/* Copy data to or from a device.  This is intended for use as an async
+   callback event.  */
+
+static void
+copy_data (void *data_)
+{
+  struct copy_data *data = (struct copy_data *)data_;
+  GCN_DEBUG ("Async thread %d:%d: Copying %zu bytes from (%p) to (%p)\n",
+     data->aq->agent->device_id, data->aq->id, data->len, data->src,
+     data->dst);
+  hsa_fns.hsa_memory_copy_fn (data->dst, data->src, data->len);
+  if (data->free_src)
+    free ((void *) data->src);
+  free (data);
+}
+
+/* Free device data.  This is intended for use as an async callback event.  */
+
+static void
+gomp_offload_free (void *ptr)
+{
+  GCN_DEBUG ("Async thread ?:?: Freeing %p\n", ptr);
+  GOMP_OFFLOAD_free (0, ptr);
+}
+
+/* Request an asynchronous data copy, to or from a device, on a given queue.
+   The event will be registered as a callback.  If FREE_SRC is true
+   then the source data will be freed following the copy.  */
+
+static void
+queue_push_copy (struct goacc_asyncqueue *aq, void *dst, const void *src,
+ size_t len, bool free_src)
+{
+  if (DEBUG_QUEUES)
+    GCN_DEBUG ("queue_push_copy %d:%d: %zu bytes from (%p) to (%p)\n",
+       aq->agent->device_id, aq->id, len, src, dst);
+  struct copy_data *data
+    = (struct copy_data *)GOMP_PLUGIN_malloc (sizeof (struct copy_data));
+  data->dst = dst;
+  data->src = src;
+  data->len = len;
+  data->free_src = free_src;
+  data->aq = aq;
+  queue_push_callback (aq, copy_data, data);
+}
+
+/* Return true if the given queue is currently empty.  */
+
+static int
+queue_empty (struct goacc_asyncqueue *aq)
+{
+  pthread_mutex_lock (&aq->mutex);
+  int res = aq->queue_n == 0 ? 1 : 0;
+  pthread_mutex_unlock (&aq->mutex);
+
+  return res;
+}
+
+/* Wait for a given queue to become empty.  This implements an OpenACC wait
+   directive.  */
+
+static void
+wait_queue (struct goacc_asyncqueue *aq)
+{
+  if (DRAIN_QUEUE_SYNCHRONOUS_P)
+    {
+      drain_queue_synchronous (aq);
+      return;
+    }
+
+  pthread_mutex_lock (&aq->mutex);
+
+  while (aq->queue_n > 0)
+    {
+      if (DEBUG_THREAD_SLEEP)
+ GCN_DEBUG ("waiting for thread %d:%d, putting thread to sleep\n",
+   aq->agent->device_id, aq->id);
+      pthread_cond_wait (&aq->queue_cond_out, &aq->mutex);
+      if (DEBUG_THREAD_SLEEP)
+ GCN_DEBUG ("thread %d:%d woke up.  Rechecking\n", aq->agent->device_id,
+   aq->id);
+    }
+
+  pthread_mutex_unlock (&aq->mutex);
+  GCN_DEBUG ("waiting for thread %d:%d, done\n", aq->agent->device_id, aq->id);
+}
+
+/* }}}  */
+/* {{{ OpenACC support  */
+
+/* Execute an OpenACC kernel, synchronously or asynchronously.  */
+
+static void
+gcn_exec (struct kernel_info *kernel, size_t mapnum, void **hostaddrs,
+  void **devaddrs, unsigned *dims, void *targ_mem_desc, bool async,
+  struct goacc_asyncqueue *aq)
+{
+  if (!GOMP_OFFLOAD_can_run (kernel))
+    GOMP_PLUGIN_fatal ("OpenACC host fallback unimplemented.");
+
+  /* If we get here then this must be an OpenACC kernel.  */
+  kernel->kind = KIND_OPENACC;
+
+  /* devaddrs must be double-indirect on the target.  */
+  void **ind_da = alloc_by_agent (kernel->agent, sizeof (void*) * mapnum);
+  for (size_t i = 0; i < mapnum; i++)
+    hsa_fns.hsa_memory_copy_fn (&ind_da[i],
+ devaddrs[i] ? &devaddrs[i] : &hostaddrs[i],
+ sizeof (void *));
+
+  struct hsa_kernel_description *hsa_kernel_desc = NULL;
+  for (unsigned i = 0; i < kernel->module->image_desc->kernel_count; i++)
+    {
+      struct hsa_kernel_description *d
+ = &kernel->module->image_desc->kernel_infos[i];
+      if (d->name == kernel->name)
+ {
+  hsa_kernel_desc = d;
+  break;
+ }
+    }
+
+  /* We may have statically-determined dimensions in
+     hsa_kernel_desc->oacc_dims[] or dimensions passed to this offload kernel
+     invocation at runtime in dims[].  We allow static dimensions to take
+     priority over dynamic dimensions when present (non-zero).  */
+  if (hsa_kernel_desc->oacc_dims[0] > 0)
+    dims[0] = hsa_kernel_desc->oacc_dims[0];
+  if (hsa_kernel_desc->oacc_dims[1] > 0)
+    dims[1] = hsa_kernel_desc->oacc_dims[1];
+  if (hsa_kernel_desc->oacc_dims[2] > 0)
+    dims[2] = hsa_kernel_desc->oacc_dims[2];
+
+  /* If any of the OpenACC dimensions remain 0 then we get to pick a number.
+     There isn't really a correct answer for this without a clue about the
+     problem size, so let's do a reasonable number of single-worker gangs.
+     64 gangs matches a typical Fiji device.  */
+
+  /* NOTE: Until support for middle-end worker partitioning is merged, use 1
+     for the default number of workers.  */
+  if (dims[0] == 0) dims[0] = get_cu_count (kernel->agent); /* Gangs.  */
+  if (dims[1] == 0) dims[1] = 1;  /* Workers.  */
+
+  /* The incoming dimensions are expressed in terms of gangs, workers, and
+     vectors.  The HSA dimensions are expressed in terms of "work-items",
+     which means multiples of vector lanes.
+
+     The "grid size" specifies the size of the problem space, and the
+     "work-group size" specifies how much of that we want a single compute
+     unit to chew on at once.
+
+     The three dimensions do not really correspond to hardware, but the
+     important thing is that the HSA runtime will launch as many
+     work-groups as it takes to process the entire grid, and each
+     work-group will contain as many wave-fronts as it takes to process
+     the work-items in that group.
+
+     Essentially, as long as we set the Y dimension to 64 (the number of
+     vector lanes in hardware), and the Z group size to the maximum (16),
+     then we will get the gangs (X) and workers (Z) launched as we expect.
+
+     The reason for the apparent reversal of vector and worker dimension
+     order is to do with the way the run-time distributes work-items across
+     v1 and v2.  */
+  struct GOMP_kernel_launch_attributes kla =
+    {3,
+     /* Grid size.  */
+     {dims[0], 64, dims[1]},
+     /* Work-group size.  */
+     {1,       64, 16}
+    };
+
+  struct goacc_thread *thr = GOMP_PLUGIN_goacc_thread ();
+  acc_prof_info *prof_info = thr->prof_info;
+  acc_event_info enqueue_launch_event_info;
+  acc_api_info *api_info = thr->api_info;
+  bool profiling_dispatch_p = __builtin_expect (prof_info != NULL, false);
+  if (profiling_dispatch_p)
+    {
+      prof_info->event_type = acc_ev_enqueue_launch_start;
+
+      enqueue_launch_event_info.launch_event.event_type
+ = prof_info->event_type;
+      enqueue_launch_event_info.launch_event.valid_bytes
+ = _ACC_LAUNCH_EVENT_INFO_VALID_BYTES;
+      enqueue_launch_event_info.launch_event.parent_construct
+ = acc_construct_parallel;
+      enqueue_launch_event_info.launch_event.implicit = 1;
+      enqueue_launch_event_info.launch_event.tool_info = NULL;
+      enqueue_launch_event_info.launch_event.kernel_name
+ = (char *) kernel->name;
+      enqueue_launch_event_info.launch_event.num_gangs = kla.gdims[0];
+      enqueue_launch_event_info.launch_event.num_workers = kla.gdims[2];
+      enqueue_launch_event_info.launch_event.vector_length = kla.gdims[1];
+
+      api_info->device_api = acc_device_api_other;
+
+      GOMP_PLUGIN_goacc_profiling_dispatch (prof_info,
+ &enqueue_launch_event_info, api_info);
+    }
+
+  if (!async)
+    {
+      run_kernel (kernel, ind_da, &kla, NULL, false);
+      gomp_offload_free (ind_da);
+    }
+  else
+    {
+      queue_push_launch (aq, kernel, ind_da, &kla);
+      if (DEBUG_QUEUES)
+ GCN_DEBUG ("queue_push_callback %d:%d gomp_offload_free, %p\n",
+   aq->agent->device_id, aq->id, ind_da);
+      queue_push_callback (aq, gomp_offload_free, ind_da);
+    }
+
+  if (profiling_dispatch_p)
+    {
+      prof_info->event_type = acc_ev_enqueue_launch_end;
+      enqueue_launch_event_info.launch_event.event_type = prof_info->event_type;
+      GOMP_PLUGIN_goacc_profiling_dispatch (prof_info,
+    &enqueue_launch_event_info,
+    api_info);
+    }
+}
+
+/* }}}  */
+/* {{{ Generic Plugin API  */
+
+/* Return the name of the accelerator, which is "gcn".  */
+
+const char *
+GOMP_OFFLOAD_get_name (void)
+{
+  return "gcn";
+}
+
+/* Return the specific capabilities the HSA accelerator have.  */
+
+unsigned int
+GOMP_OFFLOAD_get_caps (void)
+{
+  /* FIXME: Enable shared memory for APU, but not discrete GPU.  */
+  return /*GOMP_OFFLOAD_CAP_SHARED_MEM |*/ GOMP_OFFLOAD_CAP_OPENMP_400
+    | GOMP_OFFLOAD_CAP_OPENACC_200;
+}
+
+/* Identify as GCN accelerator.  */
+
+int
+GOMP_OFFLOAD_get_type (void)
+{
+  return OFFLOAD_TARGET_TYPE_GCN;
+}
+
+/* Return the libgomp version number we're compatible with.  There is
+   no requirement for cross-version compatibility.  */
+
+unsigned
+GOMP_OFFLOAD_version (void)
+{
+  return GOMP_VERSION;
+}
+
+/* Return the number of GCN devices on the system.  */
+
+int
+GOMP_OFFLOAD_get_num_devices (void)
+{
+  if (!init_hsa_context ())
+    return 0;
+  return hsa_context.agent_count;
+}
+
+/* Initialize device (agent) number N so that it can be used for computation.
+   Return TRUE on success.  */
+
+bool
+GOMP_OFFLOAD_init_device (int n)
+{
+  if (!init_hsa_context ())
+    return false;
+  if (n >= hsa_context.agent_count)
+    {
+      GOMP_PLUGIN_error ("Request to initialize non-existent GCN device %i", n);
+      return false;
+    }
+  struct agent_info *agent = &hsa_context.agents[n];
+
+  if (agent->initialized)
+    return true;
+
+  agent->device_id = n;
+
+  if (pthread_rwlock_init (&agent->module_rwlock, NULL))
+    {
+      GOMP_PLUGIN_error ("Failed to initialize a GCN agent rwlock");
+      return false;
+    }
+  if (pthread_mutex_init (&agent->prog_mutex, NULL))
+    {
+      GOMP_PLUGIN_error ("Failed to initialize a GCN agent program mutex");
+      return false;
+    }
+  if (pthread_mutex_init (&agent->async_queues_mutex, NULL))
+    {
+      GOMP_PLUGIN_error ("Failed to initialize a GCN agent queue mutex");
+      return false;
+    }
+  if (pthread_mutex_init (&agent->team_arena_write_lock, NULL))
+    {
+      GOMP_PLUGIN_error ("Failed to initialize a GCN team arena write mutex");
+      return false;
+    }
+  agent->async_queues = NULL;
+  agent->omp_async_queue = NULL;
+  agent->team_arena_list = NULL;
+
+  uint32_t queue_size;
+  hsa_status_t status;
+  status = hsa_fns.hsa_agent_get_info_fn (agent->id,
+  HSA_AGENT_INFO_QUEUE_MAX_SIZE,
+  &queue_size);
+  if (status != HSA_STATUS_SUCCESS)
+    return hsa_error ("Error requesting maximum queue size of the GCN agent",
+      status);
+
+  char buf[64];
+  status = hsa_fns.hsa_agent_get_info_fn (agent->id, HSA_AGENT_INFO_NAME,
+  &buf);
+  if (status != HSA_STATUS_SUCCESS)
+    return hsa_error ("Error querying the name of the agent", status);
+  agent->gfx900_p = (strncmp (buf, "gfx900", 6) == 0);
+
+  status = hsa_fns.hsa_queue_create_fn (agent->id, queue_size,
+ HSA_QUEUE_TYPE_MULTI,
+ hsa_queue_callback, NULL, UINT32_MAX,
+ UINT32_MAX, &agent->sync_queue);
+  if (status != HSA_STATUS_SUCCESS)
+    return hsa_error ("Error creating command queue", status);
+
+  agent->kernarg_region.handle = (uint64_t) -1;
+  status = hsa_fns.hsa_agent_iterate_regions_fn (agent->id,
+ get_kernarg_memory_region,
+ &agent->kernarg_region);
+  if (agent->kernarg_region.handle == (uint64_t) -1)
+    {
+      GOMP_PLUGIN_error ("Could not find suitable memory region for kernel "
+ "arguments");
+      return false;
+    }
+  GCN_DEBUG ("Selected kernel arguments memory region:\n");
+  dump_hsa_region (agent->kernarg_region, NULL);
+
+  agent->data_region.handle = (uint64_t) -1;
+  status = hsa_fns.hsa_agent_iterate_regions_fn (agent->id,
+ get_data_memory_region,
+ &agent->data_region);
+  if (agent->data_region.handle == (uint64_t) -1)
+    {
+      GOMP_PLUGIN_error ("Could not find suitable memory region for device "
+ "data");
+      return false;
+    }
+  GCN_DEBUG ("Selected device data memory region:\n");
+  dump_hsa_region (agent->data_region, NULL);
+
+  GCN_DEBUG ("GCN agent %d initialized\n", n);
+
+  agent->initialized = true;
+  return true;
+}
+
+/* Load GCN object-code module described by struct gcn_image_desc in
+   TARGET_DATA and return references to kernel descriptors in TARGET_TABLE.
+   If there are any constructors then run them.  */
+
+int
+GOMP_OFFLOAD_load_image (int ord, unsigned version, const void *target_data,
+ struct addr_pair **target_table)
+{
+  if (GOMP_VERSION_DEV (version) != GOMP_VERSION_GCN)
+    {
+      GOMP_PLUGIN_error ("Offload data incompatible with GCN plugin"
+ " (expected %u, received %u)",
+ GOMP_VERSION_GCN, GOMP_VERSION_DEV (version));
+      return -1;
+    }
+
+  struct gcn_image_desc *image_desc = (struct gcn_image_desc *) target_data;
+  struct agent_info *agent;
+  struct addr_pair *pair;
+  struct module_info *module;
+  struct kernel_info *kernel;
+  int kernel_count = image_desc->kernel_count;
+  unsigned var_count = image_desc->global_variable_count;
+
+  agent = get_agent_info (ord);
+  if (!agent)
+    return -1;
+
+  if (pthread_rwlock_wrlock (&agent->module_rwlock))
+    {
+      GOMP_PLUGIN_error ("Unable to write-lock a GCN agent rwlock");
+      return -1;
+    }
+  if (agent->prog_finalized
+      && !destroy_hsa_program (agent))
+    return -1;
+
+  GCN_DEBUG ("Encountered %d kernels in an image\n", kernel_count);
+  GCN_DEBUG ("Encountered %u global variables in an image\n", var_count);
+  pair = GOMP_PLUGIN_malloc ((kernel_count + var_count - 2)
+     * sizeof (struct addr_pair));
+  *target_table = pair;
+  module = (struct module_info *)
+    GOMP_PLUGIN_malloc_cleared (sizeof (struct module_info)
+ + kernel_count * sizeof (struct kernel_info));
+  module->image_desc = image_desc;
+  module->kernel_count = kernel_count;
+  module->heap = NULL;
+  module->constructors_run_p = false;
+
+  kernel = &module->kernels[0];
+
+  /* Allocate memory for kernel dependencies.  */
+  for (unsigned i = 0; i < kernel_count; i++)
+    {
+      struct hsa_kernel_description *d = &image_desc->kernel_infos[i];
+      if (!init_basic_kernel_info (kernel, d, agent, module))
+ return -1;
+      if (strcmp (d->name, "_init_array") == 0)
+ module->init_array_func = kernel;
+      else if (strcmp (d->name, "_fini_array") == 0)
+        module->fini_array_func = kernel;
+      else
+ {
+  pair->start = (uintptr_t) kernel;
+  pair->end = (uintptr_t) (kernel + 1);
+  pair++;
+ }
+      kernel++;
+    }
+
+  agent->module = module;
+  if (pthread_rwlock_unlock (&agent->module_rwlock))
+    {
+      GOMP_PLUGIN_error ("Unable to unlock a GCN agent rwlock");
+      return -1;
+    }
+
+  if (!create_and_finalize_hsa_program (agent))
+    return -1;
+
+  for (unsigned i = 0; i < var_count; i++)
+    {
+      struct global_var_info *v = &image_desc->global_variables[i];
+      GCN_DEBUG ("Looking for variable %s\n", v->name);
+
+      hsa_status_t status;
+      hsa_executable_symbol_t var_symbol;
+      status = hsa_fns.hsa_executable_get_symbol_fn (agent->executable, NULL,
+     v->name, agent->id,
+     0, &var_symbol);
+
+      if (status != HSA_STATUS_SUCCESS)
+ hsa_fatal ("Could not find symbol for variable in the code object",
+   status);
+
+      uint64_t var_addr;
+      uint32_t var_size;
+      status = hsa_fns.hsa_executable_symbol_get_info_fn
+ (var_symbol, HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_ADDRESS, &var_addr);
+      if (status != HSA_STATUS_SUCCESS)
+ hsa_fatal ("Could not extract a variable from its symbol", status);
+      status = hsa_fns.hsa_executable_symbol_get_info_fn
+ (var_symbol, HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_SIZE, &var_size);
+      if (status != HSA_STATUS_SUCCESS)
+ hsa_fatal ("Could not extract a variable size from its symbol", status);
+
+      pair->start = var_addr;
+      pair->end = var_addr + var_size;
+      GCN_DEBUG ("Found variable %s at %p with size %u\n", v->name,
+ (void *)var_addr, var_size);
+      pair++;
+    }
+
+  /* Ensure that constructors are run first.  */
+  struct GOMP_kernel_launch_attributes kla =
+    { 3,
+      /* Grid size.  */
+      { 1, 64, 1 },
+      /* Work-group size.  */
+      { 1, 64, 1 }
+    };
+
+  if (module->init_array_func)
+    {
+      init_kernel (module->init_array_func);
+      run_kernel (module->init_array_func, NULL, &kla, NULL, false);
+    }
+  module->constructors_run_p = true;
+
+  /* Don't report kernels that libgomp need not know about.  */
+  if (module->init_array_func)
+    kernel_count--;
+  if (module->fini_array_func)
+    kernel_count--;
+
+  return kernel_count + var_count;
+}
+
+/* Unload GCN object-code module described by struct gcn_image_desc in
+   TARGET_DATA from agent number N.  Return TRUE on success.  */
+
+bool
+GOMP_OFFLOAD_unload_image (int n, unsigned version, const void *target_data)
+{
+  if (GOMP_VERSION_DEV (version) != GOMP_VERSION_GCN)
+    {
+      GOMP_PLUGIN_error ("Offload data incompatible with GCN plugin"
+ " (expected %u, received %u)",
+ GOMP_VERSION_GCN, GOMP_VERSION_DEV (version));
+      return false;
+    }
+
+  struct agent_info *agent;
+  agent = get_agent_info (n);
+  if (!agent)
+    return false;
+
+  if (pthread_rwlock_wrlock (&agent->module_rwlock))
+    {
+      GOMP_PLUGIN_error ("Unable to write-lock a GCN agent rwlock");
+      return false;
+    }
+
+  if (!agent->module || agent->module->image_desc != target_data)
+    {
+      GOMP_PLUGIN_error ("Attempt to unload an image that has never been "
+ "loaded before");
+      return false;
+    }
+
+  if (!destroy_module (agent->module, true))
+    return false;
+  free (agent->module);
+  agent->module = NULL;
+  if (!destroy_hsa_program (agent))
+    return false;
+  if (pthread_rwlock_unlock (&agent->module_rwlock))
+    {
+      GOMP_PLUGIN_error ("Unable to unlock a GCN agent rwlock");
+      return false;
+    }
+  return true;
+}
+
+/* Deinitialize all information and status associated with agent number N.  We
+   do not attempt any synchronization, assuming the user and libgomp will not
+   attempt deinitialization of a device that is in any way being used at the
+   same time.  Return TRUE on success.  */
+
+bool
+GOMP_OFFLOAD_fini_device (int n)
+{
+  struct agent_info *agent = get_agent_info (n);
+  if (!agent)
+    return false;
+
+  if (!agent->initialized)
+    return true;
+
+  if (agent->omp_async_queue)
+    {
+      GOMP_OFFLOAD_openacc_async_destruct (agent->omp_async_queue);
+      agent->omp_async_queue = NULL;
+    }
+
+  if (agent->module)
+    {
+      if (!destroy_module (agent->module, false))
+ return false;
+      free (agent->module);
+      agent->module = NULL;
+    }
+
+  if (!destroy_team_arenas (agent))
+    return false;
+
+  if (!destroy_hsa_program (agent))
+    return false;
+
+  hsa_status_t status = hsa_fns.hsa_queue_destroy_fn (agent->sync_queue);
+  if (status != HSA_STATUS_SUCCESS)
+    return hsa_error ("Error destroying command queue", status);
+
+  if (pthread_mutex_destroy (&agent->prog_mutex))
+    {
+      GOMP_PLUGIN_error ("Failed to destroy a GCN agent program mutex");
+      return false;
+    }
+  if (pthread_rwlock_destroy (&agent->module_rwlock))
+    {
+      GOMP_PLUGIN_error ("Failed to destroy a GCN agent rwlock");
+      return false;
+    }
+
+  if (pthread_mutex_destroy (&agent->async_queues_mutex))
+    {
+      GOMP_PLUGIN_error ("Failed to destroy a GCN agent queue mutex");
+      return false;
+    }
+  if (pthread_mutex_destroy (&agent->team_arena_write_lock))
+    {
+      GOMP_PLUGIN_error ("Failed to destroy a GCN team arena mutex");
+      return false;
+    }
+  agent->initialized = false;
+  return true;
+}
+
+/* Return true if the HSA runtime can run function FN_PTR.  */
+
+bool
+GOMP_OFFLOAD_can_run (void *fn_ptr)
+{
+  struct kernel_info *kernel = (struct kernel_info *) fn_ptr;
+
+  init_kernel (kernel);
+  if (kernel->initialization_failed)
+    goto failure;
+
+  return true;
+
+failure:
+  if (suppress_host_fallback)
+    GOMP_PLUGIN_fatal ("GCN host fallback has been suppressed");
+  GCN_WARNING ("GCN target cannot be launched, doing a host fallback\n");
+  return false;
+}
+
+/* Allocate memory on device N.  */
+
+void *
+GOMP_OFFLOAD_alloc (int n, size_t size)
+{
+  struct agent_info *agent = get_agent_info (n);
+  return alloc_by_agent (agent, size);
+}
+
+/* Free memory from device N.  */
+
+bool
+GOMP_OFFLOAD_free (int device, void *ptr)
+{
+  GCN_DEBUG ("Freeing memory on device %d\n", device);
+
+  hsa_status_t status = hsa_fns.hsa_memory_free_fn (ptr);
+  if (status != HSA_STATUS_SUCCESS)
+    {
+      hsa_error ("Could not free device memory", status);
+      return false;
+    }
+
+  struct goacc_thread *thr = GOMP_PLUGIN_goacc_thread ();
+  bool profiling_dispatch_p
+    = __builtin_expect (thr != NULL && thr->prof_info != NULL, false);
+  if (profiling_dispatch_p)
+    {
+      acc_prof_info *prof_info = thr->prof_info;
+      acc_event_info data_event_info;
+      acc_api_info *api_info = thr->api_info;
+
+      prof_info->event_type = acc_ev_free;
+
+      data_event_info.data_event.event_type = prof_info->event_type;
+      data_event_info.data_event.valid_bytes
+ = _ACC_DATA_EVENT_INFO_VALID_BYTES;
+      data_event_info.data_event.parent_construct
+ = acc_construct_parallel;
+      data_event_info.data_event.implicit = 1;
+      data_event_info.data_event.tool_info = NULL;
+      data_event_info.data_event.var_name = NULL;
+      data_event_info.data_event.bytes = 0;
+      data_event_info.data_event.host_ptr = NULL;
+      data_event_info.data_event.device_ptr = (void *) ptr;
+
+      api_info->device_api = acc_device_api_other;
+
+      GOMP_PLUGIN_goacc_profiling_dispatch (prof_info, &data_event_info,
+    api_info);
+    }
+
+  return true;
+}
+
+/* Copy data from DEVICE to host.  */
+
+bool
+GOMP_OFFLOAD_dev2host (int device, void *dst, const void *src, size_t n)
+{
+  GCN_DEBUG ("Copying %zu bytes from device %d (%p) to host (%p)\n", n, device,
+     src, dst);
+  hsa_fns.hsa_memory_copy_fn (dst, src, n);
+  return true;
+}
+
+/* Copy data from host to DEVICE.  */
+
+bool
+GOMP_OFFLOAD_host2dev (int device, void *dst, const void *src, size_t n)
+{
+  GCN_DEBUG ("Copying %zu bytes from host (%p) to device %d (%p)\n", n, src,
+     device, dst);
+  hsa_fns.hsa_memory_copy_fn (dst, src, n);
+  return true;
+}
+
+/* Copy data within DEVICE.  Do the copy asynchronously, if appropriate.  */
+
+bool
+GOMP_OFFLOAD_dev2dev (int device, void *dst, const void *src, size_t n)
+{
+  struct gcn_thread *thread_data = gcn_thread ();
+
+  if (thread_data && !async_synchronous_p (thread_data->async))
+    {
+      struct agent_info *agent = get_agent_info (device);
+      maybe_init_omp_async (agent);
+      queue_push_copy (agent->omp_async_queue, dst, src, n, false);
+      return true;
+    }
+
+  GCN_DEBUG ("Copying %zu bytes from device %d (%p) to device %d (%p)\n", n,
+     device, src, device, dst);
+  hsa_fns.hsa_memory_copy_fn (dst, src, n);
+  return true;
+}
+
+/* }}}  */
+/* {{{ OpenMP Plugin API  */
+
+/* Run a synchronous OpenMP kernel on DEVICE and pass it an array of pointers
+   in VARS as a parameter.  The kernel is identified by FN_PTR which must point
+   to a kernel_info structure, and must have previously been loaded to the
+   specified device.  */
+
+void
+GOMP_OFFLOAD_run (int device, void *fn_ptr, void *vars, void **args)
+{
+  struct agent_info *agent = get_agent_info (device);
+  struct kernel_info *kernel = (struct kernel_info *) fn_ptr;
+  struct GOMP_kernel_launch_attributes def;
+  struct GOMP_kernel_launch_attributes *kla;
+  assert (agent == kernel->agent);
+
+  /* If we get here then the kernel must be OpenMP.  */
+  kernel->kind = KIND_OPENMP;
+
+  if (!parse_target_attributes (args, &def, &kla, agent))
+    {
+      GCN_WARNING ("Will not run GCN kernel because the grid size is zero\n");
+      return;
+    }
+  run_kernel (kernel, vars, kla, NULL, false);
+}
+
+/* Run an asynchronous OpenMP kernel on DEVICE.  This is similar to
+   GOMP_OFFLOAD_run except that the launch is queued and there is a call to
+   GOMP_PLUGIN_target_task_completion when it has finished.  */
+
+void
+GOMP_OFFLOAD_async_run (int device, void *tgt_fn, void *tgt_vars,
+ void **args, void *async_data)
+{
+  GCN_DEBUG ("GOMP_OFFLOAD_async_run invoked\n");
+  struct agent_info *agent = get_agent_info (device);
+  struct kernel_info *kernel = (struct kernel_info *) tgt_fn;
+  struct GOMP_kernel_launch_attributes def;
+  struct GOMP_kernel_launch_attributes *kla;
+  assert (agent == kernel->agent);
+
+  /* If we get here then the kernel must be OpenMP.  */
+  kernel->kind = KIND_OPENMP;
+
+  if (!parse_target_attributes (args, &def, &kla, agent))
+    {
+      GCN_WARNING ("Will not run GCN kernel because the grid size is zero\n");
+      return;
+    }
+
+  maybe_init_omp_async (agent);
+  queue_push_launch (agent->omp_async_queue, kernel, tgt_vars, kla);
+  queue_push_callback (agent->omp_async_queue,
+       GOMP_PLUGIN_target_task_completion, async_data);
+}
+
+/* }}} */
+/* {{{ OpenACC Plugin API  */
+
+/* Run a synchronous OpenACC kernel.  The device number is inferred from the
+   already-loaded KERNEL.  */
+
+void
+GOMP_OFFLOAD_openacc_exec (void (*fn_ptr) (void *), size_t mapnum,
+   void **hostaddrs, void **devaddrs, unsigned *dims,
+   void *targ_mem_desc)
+{
+  struct kernel_info *kernel = (struct kernel_info *) fn_ptr;
+
+  gcn_exec (kernel, mapnum, hostaddrs, devaddrs, dims, targ_mem_desc, false,
+    NULL);
+}
+
+/* Run an asynchronous OpenACC kernel on the specified queue.  */
+
+void
+GOMP_OFFLOAD_openacc_async_exec (void (*fn_ptr) (void *), size_t mapnum,
+ void **hostaddrs, void **devaddrs,
+ unsigned *dims, void *targ_mem_desc,
+ struct goacc_asyncqueue *aq)
+{
+  struct kernel_info *kernel = (struct kernel_info *) fn_ptr;
+
+  gcn_exec (kernel, mapnum, hostaddrs, devaddrs, dims, targ_mem_desc, true,
+    aq);
+}
+
+/* Create a new asynchronous thread and queue for running future kernels.  */
+
+struct goacc_asyncqueue *
+GOMP_OFFLOAD_openacc_async_construct (int device)
+{
+  struct agent_info *agent = get_agent_info (device);
+
+  pthread_mutex_lock (&agent->async_queues_mutex);
+
+  struct goacc_asyncqueue *aq = GOMP_PLUGIN_malloc (sizeof (*aq));
+  aq->agent = get_agent_info (device);
+  aq->prev = NULL;
+  aq->next = agent->async_queues;
+  if (aq->next)
+    {
+      aq->next->prev = aq;
+      aq->id = aq->next->id + 1;
+    }
+  else
+    aq->id = 1;
+  agent->async_queues = aq;
+
+  aq->queue_first = 0;
+  aq->queue_n = 0;
+  aq->drain_queue_stop = 0;
+
+  if (pthread_mutex_init (&aq->mutex, NULL))
+    {
+      GOMP_PLUGIN_error ("Failed to initialize a GCN agent queue mutex");
+      return false;
+    }
+  if (pthread_cond_init (&aq->queue_cond_in, NULL))
+    {
+      GOMP_PLUGIN_error ("Failed to initialize a GCN agent queue cond");
+      return false;
+    }
+  if (pthread_cond_init (&aq->queue_cond_out, NULL))
+    {
+      GOMP_PLUGIN_error ("Failed to initialize a GCN agent queue cond");
+      return false;
+    }
+
+  hsa_status_t status = hsa_fns.hsa_queue_create_fn (agent->id,
+     ASYNC_QUEUE_SIZE,
+     HSA_QUEUE_TYPE_MULTI,
+     hsa_queue_callback, NULL,
+     UINT32_MAX, UINT32_MAX,
+     &aq->hsa_queue);
+  if (status != HSA_STATUS_SUCCESS)
+    hsa_fatal ("Error creating command queue", status);
+
+  int err = pthread_create (&aq->thread_drain_queue, NULL, &drain_queue, aq);
+  if (err != 0)
+    GOMP_PLUGIN_fatal ("GCN asynchronous thread creation failed: %s",
+       strerror (err));
+  GCN_DEBUG ("Async thread %d:%d: created\n", aq->agent->device_id,
+     aq->id);
+
+  pthread_mutex_unlock (&agent->async_queues_mutex);
+
+  return aq;
+}
+
+/* Destroy an exisiting asynchronous thread and queue.  Waits for any
+   currently-running task to complete, but cancels any queued tasks.  */
+
+bool
+GOMP_OFFLOAD_openacc_async_destruct (struct goacc_asyncqueue *aq)
+{
+  struct agent_info *agent = aq->agent;
+
+  finalize_async_thread (aq);
+
+  pthread_mutex_lock (&agent->async_queues_mutex);
+
+  int err;
+  if ((err = pthread_mutex_destroy (&aq->mutex)))
+    {
+      GOMP_PLUGIN_error ("Failed to destroy a GCN async queue mutex: %d", err);
+      goto fail;
+    }
+  if (pthread_cond_destroy (&aq->queue_cond_in))
+    {
+      GOMP_PLUGIN_error ("Failed to destroy a GCN async queue cond");
+      goto fail;
+    }
+  if (pthread_cond_destroy (&aq->queue_cond_out))
+    {
+      GOMP_PLUGIN_error ("Failed to destroy a GCN async queue cond");
+      goto fail;
+    }
+  hsa_status_t status = hsa_fns.hsa_queue_destroy_fn (aq->hsa_queue);
+  if (status != HSA_STATUS_SUCCESS)
+    {
+      hsa_error ("Error destroying command queue", status);
+      goto fail;
+    }
+
+  if (aq->prev)
+    aq->prev->next = aq->next;
+  if (aq->next)
+    aq->next->prev = aq->prev;
+  if (agent->async_queues == aq)
+    agent->async_queues = aq->next;
+
+  GCN_DEBUG ("Async thread %d:%d: destroyed\n", agent->device_id, aq->id);
+
+  free (aq);
+  pthread_mutex_unlock (&agent->async_queues_mutex);
+  return true;
+
+fail:
+  pthread_mutex_unlock (&agent->async_queues_mutex);
+  return false;
+}
+
+/* Return true if the specified async queue is currently empty.  */
+
+int
+GOMP_OFFLOAD_openacc_async_test (struct goacc_asyncqueue *aq)
+{
+  return queue_empty (aq);
+}
+
+/* Block until the specified queue has executed all its tasks and the
+   queue is empty.  */
+
+bool
+GOMP_OFFLOAD_openacc_async_synchronize (struct goacc_asyncqueue *aq)
+{
+  wait_queue (aq);
+  return true;
+}
+
+/* Add a serialization point across two async queues. Any new tasks added to
+   AQ2, after this call, will not run until all tasks on AQ1, at the time
+   of this call, have completed.  */
+
+bool
+GOMP_OFFLOAD_openacc_async_serialize (struct goacc_asyncqueue *aq1,
+      struct goacc_asyncqueue *aq2)
+{
+  /* For serialize, stream aq2 waits for aq1 to complete work that has been
+     scheduled to run on it up to this point.  */
+  if (aq1 != aq2)
+    {
+      struct placeholder *placeholderp = queue_push_placeholder (aq1);
+      queue_push_asyncwait (aq2, placeholderp);
+    }
+  return true;
+}
+
+/* Add an opaque callback to the given async queue.  */
+
+void
+GOMP_OFFLOAD_openacc_async_queue_callback (struct goacc_asyncqueue *aq,
+   void (*fn) (void *), void *data)
+{
+  queue_push_callback (aq, fn, data);
+}
+
+/* Queue up an asynchronous data copy from host to DEVICE.  */
+
+bool
+GOMP_OFFLOAD_openacc_async_host2dev (int device, void *dst, const void *src,
+     size_t n, struct goacc_asyncqueue *aq)
+{
+  struct agent_info *agent = get_agent_info (device);
+  assert (agent == aq->agent);
+  /* The source data does not necessarily remain live until the deferred
+     copy happens.  Taking a snapshot of the data here avoids reading
+     uninitialised data later, but means that (a) data is copied twice and
+     (b) modifications to the copied data between the "spawning" point of
+     the asynchronous kernel and when it is executed will not be seen.
+     But, that is probably correct.  */
+  void *src_copy = GOMP_PLUGIN_malloc (n);
+  memcpy (src_copy, src, n);
+  queue_push_copy (aq, dst, src_copy, n, true);
+  return true;
+}
+
+/* Queue up an asynchronous data copy from DEVICE to host.  */
+
+bool
+GOMP_OFFLOAD_openacc_async_dev2host (int device, void *dst, const void *src,
+     size_t n, struct goacc_asyncqueue *aq)
+{
+  struct agent_info *agent = get_agent_info (device);
+  assert (agent == aq->agent);
+  queue_push_copy (aq, dst, src, n, false);
+  return true;
+}
+
+/* Set up plugin-specific thread-local-data (host-side).  */
+
+void *
+GOMP_OFFLOAD_openacc_create_thread_data (int ord __attribute__((unused)))
+{
+  struct gcn_thread *thread_data
+    = GOMP_PLUGIN_malloc (sizeof (struct gcn_thread));
+
+  thread_data->async = GOMP_ASYNC_SYNC;
+
+  return (void *) thread_data;
+}
+
+/* Clean up plugin-specific thread-local-data.  */
+
+void
+GOMP_OFFLOAD_openacc_destroy_thread_data (void *data)
+{
+  free (data);
+}
+
+/* }}} */
Reply | Threaded
Open this post in threaded view
|

Re: [PATCH 1/7 libgomp,nvptx] Move generic libgomp files from nvptx to accel

Jakub Jelinek
In reply to this post by Andrew Stubbs-4
On Tue, Nov 12, 2019 at 01:29:10PM +0000, Andrew Stubbs wrote:

> 2019-11-12  Andrew Stubbs  <[hidden email]>
>
> libgomp/
> * configure.tgt (nvptx*-*-*): Add "accel" directory.
> * config/nvptx/libgomp-plugin.c: Move ...
> * config/accel/libgomp-plugin.c: ... to here.
> * config/nvptx/lock.c: Move ...
> * config/accel/lock.c: ... to here.
> * config/nvptx/mutex.c: Move ...
> * config/accel/mutex.c: ... to here.
> * config/nvptx/mutex.h: Move ...
> * config/accel/mutex.h: ... to here.
> * config/nvptx/oacc-async.c: Move ...
> * config/accel/oacc-async.c: ... to here.
> * config/nvptx/oacc-cuda.c: Move ...
> * config/accel/oacc-cuda.c: ... to here.
> * config/nvptx/oacc-host.c: Move ...
> * config/accel/oacc-host.c: ... to here.
> * config/nvptx/oacc-init.c: Move ...
> * config/accel/oacc-init.c: ... to here.
> * config/nvptx/oacc-mem.c: Move ...
> * config/accel/oacc-mem.c: ... to here.
> * config/nvptx/oacc-plugin.c: Move ...
> * config/accel/oacc-plugin.c: ... to here.
> * config/nvptx/omp-lock.h: Move ...
> * config/accel/omp-lock.h: ... to here.
> * config/nvptx/openacc.f90: Move ...
> * config/accel/openacc.f90: ... to here.
> * config/nvptx/pool.h: Move ...
> * config/accel/pool.h: ... to here.
> * config/nvptx/proc.c: Move ...
> * config/accel/proc.c: ... to here.
> * config/nvptx/ptrlock.c: Move ...
> * config/accel/ptrlock.c: ... to here.
> * config/nvptx/ptrlock.h: Move ...
> * config/accel/ptrlock.h: ... to here.
> * config/nvptx/sem.c: Move ...
> * config/accel/sem.c: ... to here.
> * config/nvptx/sem.h: Move ...
> * config/accel/sem.h: ... to here.
> * config/nvptx/thread-stacksize.h: Move ...
> * config/accel/thread-stacksize.h: ... to here.

Ok, thanks.

        Jakub

Reply | Threaded
Open this post in threaded view
|

Re: [PATCH 3/7 libgomp,nvptx] Add device number to GOMP_OFFLOAD_openacc_async_construct

Jakub Jelinek
In reply to this post by Andrew Stubbs-4
On Tue, Nov 12, 2019 at 01:29:12PM +0000, Andrew Stubbs wrote:

> 2019-11-12  Andrew Stubbs  <[hidden email]>
>
> libgomp/
> * libgomp-plugin.h (GOMP_OFFLOAD_openacc_async_construct): Add int
> parameter.
> * oacc-async.c (lookup_goacc_asyncqueue): Pass device number to the
> queue constructor.
> * oacc-host.c (host_openacc_async_construct): Add device parameter.
> * plugin/plugin-nvptx.c (GOMP_OFFLOAD_openacc_async_construct): Add
> device parameter.

LGTM.

        Jakub

Reply | Threaded
Open this post in threaded view
|

Re: [PATCH 4/7 libgomp,amdgcn] GCN libgomp port

Jakub Jelinek
In reply to this post by Andrew Stubbs-4
On Tue, Nov 12, 2019 at 01:29:13PM +0000, Andrew Stubbs wrote:
> 2019-11-12  Andrew Stubbs  <[hidden email]>
>
> include/
> * gomp-constants.h (GOMP_DEVICE_GCN): Define.
> (GOMP_VERSION_GCN): Define.

Perhaps this could be 0, but not a big deal.

> libgomp/
> * Makefile.am (libgomp_la_SOURCES): Add oacc-target.c.
> * Makefile.in: Regenerate.
> * config.h.in (PLUGIN_GCN): Add new undef.
> * config/accel/openacc.f90 (acc_device_gcn): New parameter.
> * config/gcn/affinity-fmt.c: New file.
> * config/gcn/bar.c: New file.
> * config/gcn/bar.h: New file.
> * config/gcn/doacross.h: New file.
> * config/gcn/icv-device.c: New file.
> * config/gcn/oacc-target.c: New file.
> * config/gcn/simple-bar.h: New file.
> * config/gcn/target.c: New file.
> * config/gcn/task.c: New file.
> * config/gcn/team.c: New file.
> * config/gcn/time.c: New file.
> * configure.ac: Add amdgcn*-*-*.
> * configure: Regenerate.
> * configure.tgt: Add amdgcn*-*-*.
> * libgomp-plugin.h (offload_target_type): Add OFFLOAD_TARGET_TYPE_GCN.
> * libgomp.h (gcn_thrs): Add amdgcn variant.
> (set_gcn_thrs): Likewise.
> (gomp_thread): Likewise.
> * oacc-int.h (goacc_thread): Likewise.
> * oacc-target.c: New file.
> * openacc.f90 (acc_device_gcn): New parameter.
> * openacc.h (acc_device_t): Add acc_device_gcn.
> * team.c (gomp_free_pool_helper): Add amdgcn support.

Ok, thanks.

        Jakub

Reply | Threaded
Open this post in threaded view
|

Re: [PATCH 5/7 libgomp,amdgcn] Optimize GCN OpenMP malloc performance

Jakub Jelinek
In reply to this post by Andrew Stubbs-4
On Tue, Nov 12, 2019 at 01:29:14PM +0000, Andrew Stubbs wrote:

> 2019-11-12  Andrew Stubbs  <[hidden email]>
>
> libgomp/
> * config/gcn/team.c (gomp_gcn_enter_kernel): Set up the team arena
> and use team_malloc variants.
> (gomp_gcn_exit_kernel): Use team_free.
> * libgomp.h (TEAM_ARENA_SIZE): Define.
> (TEAM_ARENA_FREE): Define.
> (TEAM_ARENA_END): Define.
> (team_malloc): New function.
> (team_malloc_cleared): New function.
> (team_free): New function.
> * team.c (gomp_new_team): Use team_malloc.
> (free_team): Use team_free.
> (gomp_free_thread): Use team_free.
> (gomp_pause_host): Use team_free.
> * work.c (gomp_init_work_share): Use team_malloc.
> (gomp_fini_work_share): Use team_free.

> +  /* Handle OOM.  */
> +  if (result + size > *(void * __lds *)TEAM_ARENA_END)
> +    {
> +      const char msg[] = "GCN team arena exhausted\n";
> +      write (2, msg, sizeof(msg)-1);
> +      /* It's better to continue with reeduced performance than abort.

s/reeduced/reduced/

Not really sure if it is a good idea to print anything, at least not when
in some debugging mode.  I mean, it is fairly easy to write code that will
trigger this.  And, what is the reason why you can't free the
gomp_malloced memory, like comparing if the team_freed pointer is in between
TEAM_ARENA_START and TEAM_ARENA_END or similar, don't do anything in that
case, otherwise use free?

> +         Beware that this won't get freed, which might cause more problems.  */
> +      result = gomp_malloc (size);
> +    }
> +  return result;
> +}
> +
> +static inline void * __attribute__((malloc)) __attribute__((optimize("-O3")))
> +team_malloc_cleared (size_t size)
> +{
> +  char *result = team_malloc (size);
> +
> +  /* Clear the allocated memory.
> +     This should vectorize.  The allocation has been rounded up to the next
> +     4-byte boundary, so this is safe.  */
> +  for (int i = 0; i<size; i+=4)
> +    *(int*)(result+i) = 0;

Formatting (spaces around <, +=, +, between int and *.  Shouldn't 4 be
sizeof (int)?  And wouldn't memset (result, 0, size); do the same job?

        Jakub

Reply | Threaded
Open this post in threaded view
|

Re: [PATCH 7/7 libgomp,amdgcn] GCN Libgomp Plugin

Jakub Jelinek
In reply to this post by Andrew Stubbs-4
On Tue, Nov 12, 2019 at 01:29:16PM +0000, Andrew Stubbs wrote:
> 2019-11-12  Andrew Stubbs  <[hidden email]>
>
> libgomp/
> * plugin/Makefrag.am: Add amdgcn plugin support.
> * plugin/configfrag.ac: Likewise.
> * plugin/plugin-gcn.c: New file.
> * configure: Regenerate.

I'm a little bit worried about the elf.h include, not all targets might have
it, but perhaps that can be resolved incrementally if somebody reports it.

Ok.

        Jakub

Reply | Threaded
Open this post in threaded view
|

Re: [PATCH 4/7 libgomp,amdgcn] GCN libgomp port

Andrew Stubbs-4
In reply to this post by Jakub Jelinek
On 12/11/2019 13:46, Jakub Jelinek wrote:
> On Tue, Nov 12, 2019 at 01:29:13PM +0000, Andrew Stubbs wrote:
>> 2019-11-12  Andrew Stubbs  <[hidden email]>
>>
>> include/
>> * gomp-constants.h (GOMP_DEVICE_GCN): Define.
>> (GOMP_VERSION_GCN): Define.
>
> Perhaps this could be 0, but not a big deal.

OG9 uses 0 and is not binary compatible; this was a deliberate bump.

>> libgomp/
>> * Makefile.am (libgomp_la_SOURCES): Add oacc-target.c.
>> * Makefile.in: Regenerate.
>> * config.h.in (PLUGIN_GCN): Add new undef.
>> * config/accel/openacc.f90 (acc_device_gcn): New parameter.
>> * config/gcn/affinity-fmt.c: New file.
>> * config/gcn/bar.c: New file.
>> * config/gcn/bar.h: New file.
>> * config/gcn/doacross.h: New file.
>> * config/gcn/icv-device.c: New file.
>> * config/gcn/oacc-target.c: New file.
>> * config/gcn/simple-bar.h: New file.
>> * config/gcn/target.c: New file.
>> * config/gcn/task.c: New file.
>> * config/gcn/team.c: New file.
>> * config/gcn/time.c: New file.
>> * configure.ac: Add amdgcn*-*-*.
>> * configure: Regenerate.
>> * configure.tgt: Add amdgcn*-*-*.
>> * libgomp-plugin.h (offload_target_type): Add OFFLOAD_TARGET_TYPE_GCN.
>> * libgomp.h (gcn_thrs): Add amdgcn variant.
>> (set_gcn_thrs): Likewise.
>> (gomp_thread): Likewise.
>> * oacc-int.h (goacc_thread): Likewise.
>> * oacc-target.c: New file.
>> * openacc.f90 (acc_device_gcn): New parameter.
>> * openacc.h (acc_device_t): Add acc_device_gcn.
>> * team.c (gomp_free_pool_helper): Add amdgcn support.
>
> Ok, thanks.

Thanks.

Andrew
Reply | Threaded
Open this post in threaded view
|

Re: [PATCH 7/7 libgomp,amdgcn] GCN Libgomp Plugin

Andrew Stubbs-4
In reply to this post by Jakub Jelinek
On 12/11/2019 14:01, Jakub Jelinek wrote:

> On Tue, Nov 12, 2019 at 01:29:16PM +0000, Andrew Stubbs wrote:
>> 2019-11-12  Andrew Stubbs  <[hidden email]>
>>
>> libgomp/
>> * plugin/Makefrag.am: Add amdgcn plugin support.
>> * plugin/configfrag.ac: Likewise.
>> * plugin/plugin-gcn.c: New file.
>> * configure: Regenerate.
>
> I'm a little bit worried about the elf.h include, not all targets might have
> it, but perhaps that can be resolved incrementally if somebody reports it.

We only support running on x86_64 hosts that have the ROCm tools
installed. Access to elf.h ought not to be a big deal.

When we move to HSACO v3 binaries then we should be able to drop the
manual relocation handling, so this will go away then.

Thanks

Andrew
Reply | Threaded
Open this post in threaded view
|

Re: [PATCH 5/7 libgomp,amdgcn] Optimize GCN OpenMP malloc performance

Andrew Stubbs-4
In reply to this post by Jakub Jelinek
On 12/11/2019 13:56, Jakub Jelinek wrote:
> s/reeduced/reduced/

Done.

> Not really sure if it is a good idea to print anything, at least not when
> in some debugging mode.  I mean, it is fairly easy to write code that will
> trigger this.  And, what is the reason why you can't free the
> gomp_malloced memory, like comparing if the team_freed pointer is in between
> TEAM_ARENA_START and TEAM_ARENA_END or similar, don't do anything in that
> case, otherwise use free?

Falling back to malloc is a big performance hit. There's a global lock
affecting all teams in all running kernels. If we're running into this
then a) I want to know about it so I can tune the arena size, and b) I
want the user to know why performance is suddenly worse.

At least for now, I want to keep the message. I've updated the comment
though.

>> +  /* Clear the allocated memory.
>> +     This should vectorize.  The allocation has been rounded up to the next
>> +     4-byte boundary, so this is safe.  */
>> +  for (int i = 0; i<size; i+=4)
>> +    *(int*)(result+i) = 0;
>
> Formatting (spaces around <, +=, +, between int and *.  Shouldn't 4 be
> sizeof (int)?  And wouldn't memset (result, 0, size); do the same job?

I wanted to ensure that the loop would vectorize inline, but I don't
think it was doing so anyway. I need to look at that, but how is this,
for now?

Andrew

Optimize GCN OpenMP malloc performance

2019-11-12  Andrew Stubbs  <[hidden email]>

        libgomp/
        * config/gcn/team.c (gomp_gcn_enter_kernel): Set up the team arena
        and use team_malloc variants.
        (gomp_gcn_exit_kernel): Use team_free.
        * libgomp.h (TEAM_ARENA_SIZE): Define.
        (TEAM_ARENA_START): Define.
        (TEAM_ARENA_FREE): Define.
        (TEAM_ARENA_END): Define.
        (team_malloc): New function.
        (team_malloc_cleared): New function.
        (team_free): New function.
        * team.c (gomp_new_team): Initialize and use team_malloc.
        (free_team): Use team_free.
        (gomp_free_thread): Use team_free.
        (gomp_pause_host): Use team_free.
        * work.c (gomp_init_work_share): Use team_malloc.
        (gomp_fini_work_share): Use team_free.

diff --git a/libgomp/config/gcn/team.c b/libgomp/config/gcn/team.c
index c566482bda2..20d419198e0 100644
--- a/libgomp/config/gcn/team.c
+++ b/libgomp/config/gcn/team.c
@@ -57,16 +57,28 @@ gomp_gcn_enter_kernel (void)
       /* Starting additional threads is not supported.  */
       gomp_global_icv.dyn_var = true;
 
+      /* Initialize the team arena for optimized memory allocation.
+         The arena has been allocated on the host side, and the address
+         passed in via the kernargs.  Each team takes a small slice of it.  */
+      register void **kernargs asm("s8");
+      void *team_arena = (kernargs[4] + TEAM_ARENA_SIZE*teamid);
+      void * __lds *arena_start = (void * __lds *)TEAM_ARENA_START;
+      void * __lds *arena_free = (void * __lds *)TEAM_ARENA_FREE;
+      void * __lds *arena_end = (void * __lds *)TEAM_ARENA_END;
+      *arena_start = team_arena;
+      *arena_free = team_arena;
+      *arena_end = team_arena + TEAM_ARENA_SIZE;
+
       /* Allocate and initialize the team-local-storage data.  */
-      struct gomp_thread *thrs = gomp_malloc_cleared (sizeof (*thrs)
+      struct gomp_thread *thrs = team_malloc_cleared (sizeof (*thrs)
       * numthreads);
       set_gcn_thrs (thrs);
 
       /* Allocate and initailize a pool of threads in the team.
          The threads are already running, of course, we just need to manage
          the communication between them.  */
-      struct gomp_thread_pool *pool = gomp_malloc (sizeof (*pool));
-      pool->threads = gomp_malloc (sizeof (void *) * numthreads);
+      struct gomp_thread_pool *pool = team_malloc (sizeof (*pool));
+      pool->threads = team_malloc (sizeof (void *) * numthreads);
       for (int tid = 0; tid < numthreads; tid++)
  pool->threads[tid] = &thrs[tid];
       pool->threads_size = numthreads;
@@ -91,7 +103,7 @@ void
 gomp_gcn_exit_kernel (void)
 {
   gomp_free_thread (gcn_thrs ());
-  free (gcn_thrs ());
+  team_free (gcn_thrs ());
 }
 
 /* This function contains the idle loop in which a thread waits
diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h
index 19e1241ee4c..bab733d2b2d 100644
--- a/libgomp/libgomp.h
+++ b/libgomp/libgomp.h
@@ -106,6 +106,69 @@ extern void gomp_aligned_free (void *);
    GCC's builtin alloca().  */
 #define gomp_alloca(x)  __builtin_alloca(x)
 
+/* Optimized allocators for team-specific data that will die with the team.  */
+
+#ifdef __AMDGCN__
+/* The arena is initialized in config/gcn/team.c.  */
+#define TEAM_ARENA_SIZE  64*1024  /* Must match the value in plugin-gcn.c.  */
+#define TEAM_ARENA_START 16  /* LDS offset of free pointer.  */
+#define TEAM_ARENA_FREE  24  /* LDS offset of free pointer.  */
+#define TEAM_ARENA_END   32  /* LDS offset of end pointer.  */
+
+static inline void * __attribute__((malloc))
+team_malloc (size_t size)
+{
+  /* 4-byte align the size.  */
+  size = (size + 3) & ~3;
+
+  /* Allocate directly from the arena.
+     The compiler does not support DS atomics, yet. */
+  void *result;
+  asm ("ds_add_rtn_u64 %0, %1, %2\n\ts_waitcnt 0"
+       : "=v"(result) : "v"(TEAM_ARENA_FREE), "v"(size), "e"(1L) : "memory");
+
+  /* Handle OOM.  */
+  if (result + size > *(void * __lds *)TEAM_ARENA_END)
+    {
+      /* While this is experimental, let's make sure we know when OOM
+ happens.  */
+      const char msg[] = "GCN team arena exhausted\n";
+      write (2, msg, sizeof(msg)-1);
+
+      /* Fall back to using the heap (slowly).  */
+      result = gomp_malloc (size);
+    }
+  return result;
+}
+
+static inline void * __attribute__((malloc))
+team_malloc_cleared (size_t size)
+{
+  char *result = team_malloc (size);
+
+  /* Clear the allocated memory.  */
+  __builtin_memset (result, 0, size);
+
+  return result;
+}
+
+static inline void
+team_free (void *ptr)
+{
+  /* The whole arena is freed when the kernel exits.
+     However, if we fell back to using heap then we should free it.
+     It would be better if this function could be a no-op, but at least
+     LDS loads are cheap.  */
+  if (ptr < *(void * __lds *)TEAM_ARENA_START
+      || ptr >= *(void * __lds *)TEAM_ARENA_END)
+    free (ptr);
+}
+#else
+#define team_malloc(...) gomp_malloc (__VA_ARGS__)
+#define team_malloc_cleared(...) gomp_malloc_cleared (__VA_ARGS__)
+#define team_free(...) free (__VA_ARGS__)
+#endif
+
 /* error.c */
 
 extern void gomp_vdebug (int, const char *, va_list);
diff --git a/libgomp/team.c b/libgomp/team.c
index b26caaaaec6..cdfb9ba6c98 100644
--- a/libgomp/team.c
+++ b/libgomp/team.c
@@ -171,7 +171,7 @@ gomp_new_team (unsigned nthreads)
     {
       size_t extra = sizeof (team->ordered_release[0])
      + sizeof (team->implicit_task[0]);
-      team = gomp_malloc (sizeof (*team) + nthreads * extra);
+      team = team_malloc (sizeof (*team) + nthreads * extra);
 
 #ifndef HAVE_SYNC_BUILTINS
       gomp_mutex_init (&team->work_share_list_free_lock);
@@ -221,7 +221,7 @@ free_team (struct gomp_team *team)
   gomp_barrier_destroy (&team->barrier);
   gomp_mutex_destroy (&team->task_lock);
   priority_queue_free (&team->task_queue);
-  free (team);
+  team_free (team);
 }
 
 static void
@@ -285,8 +285,8 @@ gomp_free_thread (void *arg __attribute__((unused)))
       if (pool->last_team)
  free_team (pool->last_team);
 #ifndef __nvptx__
-      free (pool->threads);
-      free (pool);
+      team_free (pool->threads);
+      team_free (pool);
 #endif
       thr->thread_pool = NULL;
     }
@@ -1082,8 +1082,8 @@ gomp_pause_host (void)
       if (pool->last_team)
  free_team (pool->last_team);
 #ifndef __nvptx__
-      free (pool->threads);
-      free (pool);
+      team_free (pool->threads);
+      team_free (pool);
 #endif
       thr->thread_pool = NULL;
     }
diff --git a/libgomp/work.c b/libgomp/work.c
index a589b8b5231..28bb0c11255 100644
--- a/libgomp/work.c
+++ b/libgomp/work.c
@@ -120,7 +120,7 @@ gomp_init_work_share (struct gomp_work_share *ws, size_t ordered,
       else
  ordered = nthreads * sizeof (*ws->ordered_team_ids);
       if (ordered > INLINE_ORDERED_TEAM_IDS_SIZE)
- ws->ordered_team_ids = gomp_malloc (ordered);
+ ws->ordered_team_ids = team_malloc (ordered);
       else
  ws->ordered_team_ids = ws->inline_ordered_team_ids;
       memset (ws->ordered_team_ids, '\0', ordered);
@@ -142,7 +142,7 @@ gomp_fini_work_share (struct gomp_work_share *ws)
 {
   gomp_mutex_destroy (&ws->lock);
   if (ws->ordered_team_ids != ws->inline_ordered_team_ids)
-    free (ws->ordered_team_ids);
+    team_free (ws->ordered_team_ids);
   gomp_ptrlock_destroy (&ws->next_ws);
 }
 
Reply | Threaded
Open this post in threaded view
|

Re: [PATCH 5/7 libgomp,amdgcn] Optimize GCN OpenMP malloc performance

Jakub Jelinek
On Tue, Nov 12, 2019 at 05:47:11PM +0000, Andrew Stubbs wrote:

> > Not really sure if it is a good idea to print anything, at least not when
> > in some debugging mode.  I mean, it is fairly easy to write code that will
> > trigger this.  And, what is the reason why you can't free the
> > gomp_malloced memory, like comparing if the team_freed pointer is in between
> > TEAM_ARENA_START and TEAM_ARENA_END or similar, don't do anything in that
> > case, otherwise use free?
>
> Falling back to malloc is a big performance hit. There's a global lock
> affecting all teams in all running kernels. If we're running into this then
> a) I want to know about it so I can tune the arena size, and b) I want the
> user to know why performance is suddenly worse.
>
> At least for now, I want to keep the message. I've updated the comment
> though.

At some point we'll need to pass some data from the host to offloading
target, including ICVs etc. (at least for OpenMP 5.1) and at that point we
can pass also the GOMP_DEBUG state.

> I wanted to ensure that the loop would vectorize inline, but I don't think
> it was doing so anyway. I need to look at that, but how is this, for now?

Ok.

        Jakub

Reply | Threaded
Open this post in threaded view
|

Re: [PATCH 0/7 libgomp,amdgcn] AMD GCN Offloading Support

Andrew Stubbs-4
In reply to this post by Andrew Stubbs-4
These patches are now all committed. I've adjusted the changelogs to
list all the proper authors (apologies if I missed anyone).

Thank you for the quick reviews, Jakub. :-)

Andrew

On 12/11/2019 13:29, Andrew Stubbs wrote:

> Hi all,
>
> This patch series contributes initial OpenMP and OpenACC support for AMD
> GCN GPUs.
>
> The test results are not yet perfect, but there are many more passes than
> failures, so this is a good starting point.  The rest of the issues can
> be addressed as bugs during stage 3.
>
> I have another, unfinished, patch to massage the testsuite itself.  I'll
> post this shortly, once I've finished checking the forward port is
> appropriate.
>
> This series implements only single-worker support for OpenACC.  Julian
> Brown may post the multiple-worker support soon, if it isn't too
> difficult to forward-port.  Otherwise that will have to wait for GCC 11.
>
> Andrew
>
> Andrew Stubbs (7):
>    Move generic libgomp files from nvptx to accel
>    GCN mkoffload
>    Add device number to GOMP_OFFLOAD_openacc_async_construct
>    GCN libgomp port
>    Optimize GCN OpenMP malloc performance
>    Use a single worker for OpenACC on AMD GCN
>    GCN Libgomp Plugin
>
>   gcc/config/gcn/gcn.c                          |    4 +-
>   gcc/config/gcn/gcn.opt                        |    2 +-
>   gcc/config/gcn/mkoffload.c                    |  694 +++
>   gcc/config/gcn/offload.h                      |   35 +
>   include/gomp-constants.h                      |    2 +
>   libgomp/Makefile.am                           |    2 +-
>   libgomp/Makefile.in                           |   61 +-
>   libgomp/config.h.in                           |    3 +
>   .../config/{nvptx => accel}/libgomp-plugin.c  |    0
>   libgomp/config/{nvptx => accel}/lock.c        |    0
>   libgomp/config/{nvptx => accel}/mutex.c       |    0
>   libgomp/config/{nvptx => accel}/mutex.h       |    0
>   libgomp/config/{nvptx => accel}/oacc-async.c  |    0
>   libgomp/config/{nvptx => accel}/oacc-cuda.c   |    0
>   libgomp/config/{nvptx => accel}/oacc-host.c   |    0
>   libgomp/config/{nvptx => accel}/oacc-init.c   |    0
>   libgomp/config/{nvptx => accel}/oacc-mem.c    |    0
>   libgomp/config/{nvptx => accel}/oacc-plugin.c |    0
>   libgomp/config/{nvptx => accel}/omp-lock.h    |    0
>   libgomp/config/{nvptx => accel}/openacc.f90   |    1 +
>   libgomp/config/{nvptx => accel}/pool.h        |    0
>   libgomp/config/{nvptx => accel}/proc.c        |    0
>   libgomp/config/{nvptx => accel}/ptrlock.c     |    0
>   libgomp/config/{nvptx => accel}/ptrlock.h     |    0
>   libgomp/config/{nvptx => accel}/sem.c         |    0
>   libgomp/config/{nvptx => accel}/sem.h         |    0
>   .../{nvptx => accel}/thread-stacksize.h       |    0
>   libgomp/config/gcn/affinity-fmt.c             |   51 +
>   libgomp/config/gcn/bar.c                      |  232 +
>   libgomp/config/gcn/bar.h                      |  168 +
>   libgomp/config/gcn/doacross.h                 |   58 +
>   libgomp/config/gcn/icv-device.c               |   72 +
>   libgomp/config/gcn/oacc-target.c              |   31 +
>   libgomp/config/gcn/simple-bar.h               |   61 +
>   libgomp/config/gcn/target.c                   |   67 +
>   libgomp/config/gcn/task.c                     |   39 +
>   libgomp/config/gcn/team.c                     |  212 +
>   libgomp/config/gcn/time.c                     |   52 +
>   libgomp/configure                             |   61 +-
>   libgomp/configure.ac                          |    2 +-
>   libgomp/configure.tgt                         |    6 +-
>   libgomp/libgomp-plugin.h                      |    5 +-
>   libgomp/libgomp.h                             |   74 +
>   libgomp/oacc-async.c                          |    3 +-
>   libgomp/oacc-host.c                           |    2 +-
>   libgomp/oacc-int.h                            |    9 +-
>   libgomp/oacc-target.c                         |    1 +
>   libgomp/openacc.f90                           |    1 +
>   libgomp/openacc.h                             |    1 +
>   libgomp/plugin/Makefrag.am                    |   14 +
>   libgomp/plugin/configfrag.ac                  |   35 +
>   libgomp/plugin/plugin-gcn.c                   | 3985 +++++++++++++++++
>   libgomp/plugin/plugin-nvptx.c                 |    2 +-
>   libgomp/team.c                                |   15 +-
>   libgomp/work.c                                |    4 +-
>   55 files changed, 6035 insertions(+), 32 deletions(-)
>   create mode 100644 gcc/config/gcn/mkoffload.c
>   create mode 100644 gcc/config/gcn/offload.h
>   rename libgomp/config/{nvptx => accel}/libgomp-plugin.c (100%)
>   rename libgomp/config/{nvptx => accel}/lock.c (100%)
>   rename libgomp/config/{nvptx => accel}/mutex.c (100%)
>   rename libgomp/config/{nvptx => accel}/mutex.h (100%)
>   rename libgomp/config/{nvptx => accel}/oacc-async.c (100%)
>   rename libgomp/config/{nvptx => accel}/oacc-cuda.c (100%)
>   rename libgomp/config/{nvptx => accel}/oacc-host.c (100%)
>   rename libgomp/config/{nvptx => accel}/oacc-init.c (100%)
>   rename libgomp/config/{nvptx => accel}/oacc-mem.c (100%)
>   rename libgomp/config/{nvptx => accel}/oacc-plugin.c (100%)
>   rename libgomp/config/{nvptx => accel}/omp-lock.h (100%)
>   rename libgomp/config/{nvptx => accel}/openacc.f90 (98%)
>   rename libgomp/config/{nvptx => accel}/pool.h (100%)
>   rename libgomp/config/{nvptx => accel}/proc.c (100%)
>   rename libgomp/config/{nvptx => accel}/ptrlock.c (100%)
>   rename libgomp/config/{nvptx => accel}/ptrlock.h (100%)
>   rename libgomp/config/{nvptx => accel}/sem.c (100%)
>   rename libgomp/config/{nvptx => accel}/sem.h (100%)
>   rename libgomp/config/{nvptx => accel}/thread-stacksize.h (100%)
>   create mode 100644 libgomp/config/gcn/affinity-fmt.c
>   create mode 100644 libgomp/config/gcn/bar.c
>   create mode 100644 libgomp/config/gcn/bar.h
>   create mode 100644 libgomp/config/gcn/doacross.h
>   create mode 100644 libgomp/config/gcn/icv-device.c
>   create mode 100644 libgomp/config/gcn/oacc-target.c
>   create mode 100644 libgomp/config/gcn/simple-bar.h
>   create mode 100644 libgomp/config/gcn/target.c
>   create mode 100644 libgomp/config/gcn/task.c
>   create mode 100644 libgomp/config/gcn/team.c
>   create mode 100644 libgomp/config/gcn/time.c
>   create mode 100644 libgomp/oacc-target.c
>   create mode 100644 libgomp/plugin/plugin-gcn.c
>

Reply | Threaded
Open this post in threaded view
|

Re: [PATCH 4/7 libgomp,amdgcn] GCN libgomp port

Thomas Schwinge-8
In reply to this post by Andrew Stubbs-4
Hi!

On 2019-11-12T13:29:13+0000, Andrew Stubbs <[hidden email]> wrote:
> --- a/include/gomp-constants.h
> +++ b/include/gomp-constants.h
> @@ -174,6 +174,7 @@ enum gomp_map_kind
>  #define GOMP_DEVICE_NVIDIA_PTX 5
>  #define GOMP_DEVICE_INTEL_MIC 6
>  #define GOMP_DEVICE_HSA 7
> +#define GOMP_DEVICE_GCN 8

Unless we are to expect non-AMD GCN implementations (hardware), I'd favor
if all these "GCN" things were called "AMD GCN", like done for "Nvidia
PTX", or "Intel MIC", or why shouldn't we?


> --- a/libgomp/configure.ac
> +++ b/libgomp/configure.ac
> @@ -176,7 +176,7 @@ case "$host" in
>    *-*-rtems*)
>      # RTEMS supports Pthreads, but the library is not available at GCC build time.
>      ;;
> -  nvptx*-*-*)
> +  nvptx*-*-* | amdgcn*-*-*)
>      # NVPTX does not support Pthreads, has its own code replacement.
>      libgomp_use_pthreads=no
>      # NVPTX is an accelerator-only target
I'd have sorted alphabetically, and updated the "NVPTX" comments.

> --- a/libgomp/configure.tgt
> +++ b/libgomp/configure.tgt
> @@ -164,6 +164,10 @@ case "${target}" in
>   fi
>   ;;
>  
> +  amdgcn*-*-*)
> + config_path="gcn accel"
> + ;;
> +
>    *)
>   ;;
(Again, I'd sort alphabetically, but again, I acknowledge that the list is
not sorted at present.)


> --- a/libgomp/openacc.h
> +++ b/libgomp/openacc.h
> @@ -55,6 +55,7 @@ typedef enum acc_device_t {
>    /* acc_device_host_nonshm = 3 removed.  */
>    acc_device_not_host = 4,
>    acc_device_nvidia = 5,
> +  acc_device_gcn = 8,

There is no 'acc_device_gcn' in OpenACC.  Per OpenACC 3.0, A.1.2. "AMD
GPU Targets", for example, there is 'acc_device_radeon' (and "the
case-insensitive name 'radeon' for the environment variable
'ACC_DEVICE_TYPE'").  If that is not appropriate to use (I have not read
up how exactly AMD's "GCN" and "radeon" relate to each other), we should
get that clarified in the OpenACC specification.


Grüße
 Thomas

signature.asc (671 bytes) Download Attachment
Reply | Threaded
Open this post in threaded view
|

Re: [PATCH 4/7 libgomp,amdgcn] GCN libgomp port

Julian Brown-2
On Mon, 2 Dec 2019 15:43:29 +0100
Thomas Schwinge <[hidden email]> wrote:

> > --- a/libgomp/openacc.h
> > +++ b/libgomp/openacc.h
> > @@ -55,6 +55,7 @@ typedef enum acc_device_t {
> >    /* acc_device_host_nonshm = 3 removed.  */
> >    acc_device_not_host = 4,
> >    acc_device_nvidia = 5,
> > +  acc_device_gcn = 8,  
>
> There is no 'acc_device_gcn' in OpenACC.  Per OpenACC 3.0, A.1.2. "AMD
> GPU Targets", for example, there is 'acc_device_radeon' (and "the
> case-insensitive name 'radeon' for the environment variable
> 'ACC_DEVICE_TYPE'").  If that is not appropriate to use (I have not
> read up how exactly AMD's "GCN" and "radeon" relate to each other),
> we should get that clarified in the OpenACC specification.

FWIW, I'm pretty sure there are Radeon devices that do not use the GCN
ISA. OTOH, there are also Nvidia devices that are not PTX-compatible.
No strong opinion on the acc_device_foo name from me, though (maybe
"acc_device_amdgcn"?).

Julian
12