/* Plugin for AMD GCN execution.

   Copyright (C) 2013-2026 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  */

#define _GNU_SOURCE
#include "config.h"
#include "symcat.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 <hsa_ext_amd.h>
#include <dlfcn.h>
#include <signal.h>
#include "alloc_cache.h"
#define _LIBGOMP_PLUGIN_INCLUDE 1
#include "libgomp-plugin.h"
#undef _LIBGOMP_PLUGIN_INCLUDE
#include "config/gcn/libgomp-gcn.h"  /* For struct output.  */
#include "gomp-constants.h"
#include <elf.h>
#include "oacc-plugin.h"
#include "oacc-int.h"
#include <assert.h>
#include <sys/mman.h>
#include <unistd.h>

/* Create hash-table for declare target's indirect clause on the host;
   see build-target-indirect-htab.h for details.  */
#define USE_HASHTAB_LOOKUP_FOR_INDIRECT
#ifdef USE_HASHTAB_LOOKUP_FOR_INDIRECT
static void* create_target_indirect_map (size_t *, size_t,
					 uint64_t *, uint64_t *);
#endif

/* These probably won't be in elf.h for a while.  */
#ifndef R_AMDGPU_NONE
#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 R_AMDGPU_RELATIVE64	13	/* B + A  */
#endif

#define ELFABIVERSION_AMDGPU_HSA_V6		4

#define EF_AMDGPU_GENERIC_VERSION_V		0xff000000  /* Mask.  */
#define EF_AMDGPU_GENERIC_VERSION_OFFSET	24

#define GET_GENERIC_VERSION(VAR) ((VAR & EF_AMDGPU_GENERIC_VERSION_V) \
				  >> EF_AMDGPU_GENERIC_VERSION_OFFSET)

/* 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 implementation 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);
  uint64_t (*hsa_queue_load_read_index_relaxed_fn) (const hsa_queue_t *queue);
  uint64_t (*hsa_queue_load_write_index_relaxed_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);
  hsa_status_t (*hsa_amd_memory_fill_fn)(void *ptr, uint32_t value,
					 size_t count);
  hsa_status_t (*hsa_amd_memory_lock_fn)
    (void *host_ptr, size_t size, hsa_agent_t *agents, int num_agent,
     void **agent_ptr);
  hsa_status_t (*hsa_amd_memory_unlock_fn) (void *host_ptr);
  hsa_status_t (*hsa_amd_memory_async_copy_rect_fn)
    (const hsa_pitched_ptr_t *dst, const hsa_dim3_t *dst_offset,
     const hsa_pitched_ptr_t *src, const hsa_dim3_t *src_offset,
     const hsa_dim3_t *range, hsa_agent_t copy_agent,
     hsa_amd_copy_direction_t dir, uint32_t num_dep_signals,
     const hsa_signal_t *dep_signals, hsa_signal_t completion_signal);
  hsa_status_t (*hsa_amd_svm_attributes_set_fn)
    (void* ptr, size_t size, hsa_amd_svm_attribute_pair_t* attribute_list,
     size_t attribute_count);
  hsa_status_t (*hsa_amd_svm_attributes_get_fn)
    (void* ptr, size_t size, hsa_amd_svm_attribute_pair_t* attribute_list,
     size_t attribute_count);
  hsa_status_t (*hsa_amd_pointer_info_fn)
    (const void *, hsa_amd_pointer_info_t *, void *(*)(size_t),
     uint32_t *, hsa_agent_t **); 
};

/* As an HIP runtime is dlopened, following structure defines function
   pointers utilized by the interop feature of this plugin.
   Add sufficient type declarations to get this work.  */

typedef int hipError_t;  /* Actually an enum; 0 == success. */
typedef void* hipCtx_t;
struct hipStream_s;
typedef struct hipStream_s* hipStream_t;

struct hip_runtime_fn_info
{
  hipError_t (*hipStreamCreate_fn) (hipStream_t *);
  hipError_t (*hipStreamDestroy_fn) (hipStream_t);
  hipError_t (*hipStreamSynchronize_fn) (hipStream_t);
  hipError_t (*hipCtxGetCurrent_fn) (hipCtx_t *ctx);
  hipError_t (*hipSetDevice_fn) (int deviceId);
  hipError_t (*hipGetDevice_fn) (int *deviceId);
};

/* Structure describing the run-time and grid properties of an HSA kernel
   launch.  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.
   Redundant with parts of hsa_kernel_dispatch_packet_t.  It is maintained
   separately because the contents of the aforementioned packet become
   unspecified after dispatch, so, we can't re-read back pointers we wrote into
   the dispatch packet in order to clean them up.  */

struct kernel_dispatch
{
  struct agent_info *agent;
  /* Pointer to a command queue associated with a kernel dispatch agent.  */
  void *queue;
  /* Kernel object.  */
  uint64_t object;
  /* Synchronization signal used for dispatch synchronization.  */
  hsa_signal_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 {
  struct kernargs_abi abi;

  /* Output data.  */
  struct output output_data;

  /* Target variable table.  Size determined by gomp_map_vars.  See
     GOMP_OFFLOAD_session_allocate_target_var_table.  */
  _Alignas (__BIGGEST_ALIGNMENT__) void *target_variable_table[];
};

/* A queue entry for a future asynchronous launch.  */

struct kernel_launch
{
  struct gomp_offload_session *session;
  struct kernel_info *kernel;
  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;
  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.  */
  int sgpr_count;
  int vpgr_count;
};

/* 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 ind_func_count;
  const unsigned global_variable_count;
};

/* Enum values corresponding to the the ELF architecture codes.
   Only 'special' values are actually referenced in this file, but having them
   all may aid debugging.  */

typedef enum {
  EF_AMDGPU_MACH_UNSUPPORTED = -1,
#define GCN_DEVICE(name, NAME, ELF, ...) \
  EF_AMDGPU_MACH_AMDGCN_ ## NAME = ELF,
#include "../../gcc/config/gcn/gcn-devices.def"
} EF_AMDGPU_MACH;

const static int EF_AMDGPU_MACH_MASK = 0x000000ff;
typedef EF_AMDGPU_MACH gcn_isa;

/* 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;

  /* The instruction set architecture of the device. */
  gcn_isa device_isa;
  /* Name of the agent. */
  char name[64];
  /* Name of the vendor of the agent. */
  char vendor_name[64];
  /* 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;

  /* A stack of allocations in kernarg_region of (sizeof (struct kernargs))
     size each, used for ammortizing kernel argument allocation cost.  */
  struct alloc_cache kernarg_cache;

  /* The HSA memory region from which to allocate device data.  */
  hsa_region_t data_region;

  /* Allocated ephemeral memories (team arena and stack space).  */
  struct ephemeral_memories_list *ephemeral_memories_list;
  pthread_mutex_t ephemeral_memories_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;
  /* Information provided by mkoffload associated with the kernel.  */
  struct hsa_kernel_description *description;
  /* 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 used by OpenMP, as a means to optimize per-team malloc,
   and for host-accessible stack space.  */

struct ephemeral_memories_list
{
  struct ephemeral_memories_list *next;

  /* The size is determined by the number of teams and threads.  */
  size_t size;
  /* The device address allocated memory.  */
  void *address;
  /* A flag to prevent two asynchronous kernels trying to use the same memory.
     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;
  /* Driver version string. */
  char driver_version_s[30];
};

/* }}}  */
/* {{{ 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;

/* HIP runtime functions that are initialized in init_hip_runtime_functions.  */
static struct hip_runtime_fn_info hip_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;

/* Ephemeral memory sizes for each kernel launch.  */

static int team_arena_size = DEFAULT_TEAM_ARENA_SIZE;
static int stack_size = DEFAULT_GCN_STACK_SIZE;
static int lowlat_size = -1;

/* 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 whether to prevent merging the kernel arguments with the
   target variable table, i.e. whether to always fail
   GOMP_OFFLOAD_session_allocate_target_var_table.  Set from the
   GCN_INHIBIT_KERNARGS_TVT_MERGE env var.  */

static bool inhibit_kernargs_tvt_merge;

/* Flag to decide if the runtime should suppress a possible fallback to host
   execution.  */

static bool suppress_host_fallback;

/* Flag to locate HSA and HIP runtime shared libraries that are dlopened
   by this plug-in.  */

static const char *hsa_runtime_lib;
static const char *hip_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");

  bool svm_supported;
  status = hsa_fns.hsa_system_get_info_fn
    (HSA_AMD_SYSTEM_INFO_SVM_SUPPORTED, &svm_supported);
  if (status == HSA_STATUS_SUCCESS)
    GCN_DEBUG ("HSA_AMD_SYSTEM_INFO_SVM_SUPPORTED: %s\n",
	       (svm_supported ? "TRUE" : "FALSE"));
  else
    GCN_WARNING ("HSA_AMD_SYSTEM_INFO_SVM_SUPPORTED: FAILED\n");

  bool svm_accessible;
  status = hsa_fns.hsa_system_get_info_fn
    (HSA_AMD_SYSTEM_INFO_SVM_ACCESSIBLE_BY_DEFAULT, &svm_accessible);
  if (status == HSA_STATUS_SUCCESS)
    GCN_DEBUG ("HSA_AMD_SYSTEM_INFO_SVM_ACCESSIBLE_BY_DEFAULT: %s\n",
	       (svm_accessible ? "TRUE" : "FALSE"));
  else
    GCN_WARNING ("HSA_AMD_SYSTEM_INFO_SVM_ACCESSIBLE_BY_DEFAULT: 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 inline void
print_kernel_dispatch (struct kernel_dispatch *dispatch, unsigned indent,
		       struct kernargs *kernargs)
{
  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->abi.heap_ptr);
  fprintf (stderr, "%*sarena address: %p (%d bytes per workgroup)\n", indent,
	   "", (void*)kernargs->abi.arena_ptr,
	   kernargs->abi.arena_size_per_team);
  fprintf (stderr, "%*sstack address: %p (%d bytes per wavefront)\n", indent,
	   "", (void*)kernargs->abi.stack_ptr,
	   kernargs->abi.stack_size_per_thread);
  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 (low-latency pool)\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 = "libhsa-runtime64.so.1";

  hip_runtime_lib = secure_getenv ("HIP_RUNTIME_LIB");
  if (hip_runtime_lib == NULL)
    hip_runtime_lib = "libamdhip64.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;
    }

  const char *arena = secure_getenv ("GCN_TEAM_ARENA_SIZE");
  if (arena)
    {
      int tmp = atoi (arena);
      if (tmp)
	team_arena_size = tmp;;
    }

  const char *stack = secure_getenv ("GCN_STACK_SIZE");
  if (stack)
    {
      int tmp = atoi (stack);
      if (tmp)
	stack_size = tmp;;
    }

  const char *lowlat = secure_getenv ("GOMP_GCN_LOWLAT_POOL");
  if (lowlat)
    lowlat_size = atoi (lowlat);

  inhibit_kernargs_tvt_merge
    = (bool) secure_getenv ("GCN_INHIBIT_KERNARGS_TVT_MERGE");
}

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

/* 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 intelligent 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;
}

/* This sets the maximum number of teams to twice the number of GPU Compute
   Units to avoid memory waste and corresponding memory access faults.  */

static int
limit_teams (int teams, struct agent_info *agent)
{
  int max_teams = 2 * get_cu_count (agent);
  if (teams > max_teams)
    teams = max_teams;
  return teams;
}

/* 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_MASK)
	       == GOMP_TARGET_ARG_DEVICE_ALL)
	{
	  gcn_dims_found = true;
	  switch (id & GOMP_TARGET_ARG_ID_MASK)
	    {
	    case GOMP_TARGET_ARG_NUM_TEAMS:
	      gcn_teams = limit_teams (val, agent);
	      break;
	    case GOMP_TARGET_ARG_THREAD_LIMIT:
	      gcn_threads = limit_worker_threads (val);
	      break;
	    default:
	      ;
	    }
	}
    }

  if (gcn_dims_found)
    {
      bool gfx900_workaround_p = false;

      if (agent->device_isa == EF_AMDGPU_MACH_AMDGCN_GFX900
	  && gcn_threads == 0 && override_z_dim == 0)
	{
	  gfx900_workaround_p = true;
	  GCN_WARNING ("VEGA BUG WORKAROUND: reducing default number of "
		       "threads to at most 4 per team.\n");
	  GCN_WARNING (" - If this is not a Vega 10 device, please use "
		       "GCN_NUM_THREADS=16\n");
	}

      /* Ideally, when a dimension isn't explicitly specified, we should
	 tune it to run 40 (or 32?) threads per CU with no threads getting queued.
	 In practice, we tune for peak performance on BabelStream, which
	 for OpenACC is currently 32 threads per CU.  */
      def->ndim = 3;
      if (gcn_teams <= 0 && gcn_threads <= 0)
	{
	  /* Set up a reasonable number of teams and threads.  */
	  gcn_threads = gfx900_workaround_p ? 4 : 16; // 8;
	  def->gdims[0] = get_cu_count (agent); // * (40 / gcn_threads);
	  def->gdims[2] = gcn_threads;
	}
      else if (gcn_teams <= 0 && gcn_threads > 0)
	{
	  /* Auto-scale the number of teams with the number of threads.  */
	  def->gdims[0] = get_cu_count (agent); // * (40 / gcn_threads);
	  def->gdims[2] = gcn_threads;
	}
      else if (gcn_teams > 0 && gcn_threads <= 0)
	{
	  int max_threads = gfx900_workaround_p ? 4 : 16;

	  /* Auto-scale the number of threads with the number of teams.  */
	  def->gdims[0] = gcn_teams;
	  def->gdims[2] = 16; // get_cu_count (agent) * 40 / gcn_teams;
	  if (def->gdims[2] == 0)
	    def->gdims[2] = 1;
	  else if (def->gdims[2] > max_threads)
	    def->gdims[2] = max_threads;
	}
      else
	{
	  def->gdims[0] = gcn_teams;
	  def->gdims[2] = gcn_threads;
	}
      def->gdims[1] = 64; /* Each thread is 64 work items wide.  */
      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) \
    GOMP_PLUGIN_fatal ("'%s' is missing '%s'", hsa_runtime_lib, #function);
#define DLSYM_OPT_FN(function) \
  hsa_fns.function##_fn = dlsym (handle, #function);

  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_queue_load_read_index_relaxed)
  DLSYM_FN (hsa_queue_load_write_index_relaxed)
  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)
  DLSYM_OPT_FN (hsa_amd_memory_fill)
  DLSYM_OPT_FN (hsa_amd_memory_lock)
  DLSYM_OPT_FN (hsa_amd_memory_unlock)
  DLSYM_OPT_FN (hsa_amd_memory_async_copy_rect)
  DLSYM_OPT_FN (hsa_amd_svm_attributes_set)
  DLSYM_OPT_FN (hsa_amd_svm_attributes_get)
  DLSYM_OPT_FN (hsa_amd_pointer_info)
  return true;
#undef DLSYM_OPT_FN
#undef DLSYM_FN
}

static gcn_isa isa_code (const char *isa);

/* 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:
      {
	char name[64];
	hsa_status_t status
	  = hsa_fns.hsa_agent_get_info_fn (agent, HSA_AGENT_INFO_NAME, name);
	if (status != HSA_STATUS_SUCCESS
	    || isa_code (name) == EF_AMDGPU_MACH_UNSUPPORTED)
	  {
	    GCN_DEBUG ("Ignoring unsupported agent '%s'\n",
		       status == HSA_STATUS_SUCCESS ? name : "invalid");
	    return false;
	  }
      }
      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.
   If !PROBE: returns TRUE on success.
   If PROBE: returns TRUE on success or if the plugin/device shall be silently
   ignored, and otherwise emits an error and returns FALSE.  */

static bool
init_hsa_context (bool probe)
{
  hsa_status_t status;
  int agent_index = 0;

  if (hsa_context.initialized)
    return true;
  init_environment_variables ();
  if (!init_hsa_runtime_functions ())
    {
      const char *msg = "Run-time could not be dynamically opened";
      if (suppress_host_fallback)
	GOMP_PLUGIN_fatal ("%s\n", msg);
      else
	GCN_WARNING ("%s\n", msg);
      return probe ? true : 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 (status != HSA_STATUS_SUCCESS)
    return hsa_error ("Scanning compute agents failed", status);
  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");
    }

  uint16_t minor, major;
  status = hsa_fns.hsa_system_get_info_fn (HSA_SYSTEM_INFO_VERSION_MINOR,
					   &minor);
  if (status != HSA_STATUS_SUCCESS)
    GOMP_PLUGIN_error ("Failed to obtain HSA runtime minor version");
  status = hsa_fns.hsa_system_get_info_fn (HSA_SYSTEM_INFO_VERSION_MAJOR,
					   &major);
  if (status != HSA_STATUS_SUCCESS)
    GOMP_PLUGIN_error ("Failed to obtain HSA runtime major version");

  size_t len = sizeof hsa_context.driver_version_s;
  int printed = snprintf (hsa_context.driver_version_s, len,
			  "HSA Runtime %hu.%hu", (unsigned short int)major,
			  (unsigned short int)minor);
  if (printed >= len)
    GCN_WARNING ("HSA runtime version string was truncated."
		 "Version %hu.%hu is too long.", (unsigned short int)major,
		 (unsigned short int)minor);

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

static int
elf_gcn_isa_field (Elf64_Ehdr *image)
{
  return image->e_flags & EF_AMDGPU_MACH_MASK;
}

static int
elf_gcn_isa_is_generic (Elf64_Ehdr *image)
{
  return (image->e_ident[8] == ELFABIVERSION_AMDGPU_HSA_V6
	  && GET_GENERIC_VERSION (image->e_flags));
}

/* Returns the name that the HSA runtime uses for the ISA or NULL if we do not
   support the ISA. */

static const char*
isa_name (int isa) {
  switch(isa)
    {
#define GCN_DEVICE(name, NAME, ELF, ...) \
    case ELF: return #name;
#include "../../gcc/config/gcn/gcn-devices.def"
    }
  return NULL;
}

/* Returns the code which is used in the GCN object code to identify the ISA with
   the given name (as used by the HSA runtime).  */

static gcn_isa
isa_code(const char *isa) {
#define GCN_DEVICE(name, NAME, ELF, ...) \
  if (!strcmp (isa, #name)) return ELF;
#include "../../gcc/config/gcn/gcn-devices.def"

  return EF_AMDGPU_MACH_UNSUPPORTED;
}

/* Returns the code which is used in the GCN object code to identify the
   generic ISA that corresponds to a specific ISA.  */

static gcn_isa
generic_isa_code (int isa) {
  switch(isa)
    {
#define EF_AMDGPU_MACH_AMDGCN_NONE 0
#define GCN_DEVICE(name, NAME, ELF, GCCISA, XNACK, SRAM, WAVE64, CUMODE, \
		   VGPRS, CO, ARCH, GENERIC_ISA, ...) \
    case ELF: return EF_AMDGPU_MACH_AMDGCN_ ## GENERIC_ISA;
#include "../../gcc/config/gcn/gcn-devices.def"
    }
  return 0;
}

/* CDNA2 devices have twice as many VGPRs compared to older devices.  */

static int
max_isa_vgprs (int isa)
{
  switch (isa)
    {
#define GCN_DEVICE(name, NAME, ELF, ISA, XNACK, SRAM, WAVE64, CU, \
		   MAX_ISA_VGPRS, ...) \
    case ELF: return MAX_ISA_VGPRS;
#include "../../gcc/config/gcn/gcn-devices.def"
    default:
      GOMP_PLUGIN_fatal ("unhandled ISA in max_isa_vgprs");
    }
}

/* }}}  */
/* {{{ Run  */
struct gomp_offload_session
{
  /* Pointer to a memory space used for kernel arguments passing, wrapped in a
     node from the agent kernel argument cache.  */
  struct alloc_cache_node *kernarg_cache_node;
  /* Pointer to the actual target variable table.  */
  void **target_var_table;
  /* Device executing the kernel for this offload session.  */
  struct agent_info *agent;
};
GOMP_OFFLOAD_session_boilerplate ();

/* Prepare SESSION for use by AGENT.  */
static inline void
init_session (struct gomp_offload_session *session, struct agent_info *agent)
{
  assert (agent);
  *session = (struct gomp_offload_session) {
    .kernarg_cache_node = NULL,
    .target_var_table = NULL,
    .agent = agent,
  };
}

void
GOMP_OFFLOAD_session_start (struct gomp_offload_session *session, int device)
{
  GCN_DEBUG ("Starting session %p\n", session);
  assert ((((uintptr_t) session) % __BIGGEST_ALIGNMENT__) == 0);
  init_session (session, get_agent_info (device));
}

/* Release resources held by SESSION (but not SESSION itself).  */
static inline void
release_session (struct gomp_offload_session *session)
{
  release_alloc_cache_node (session->kernarg_cache_node);
}

/* Get new kernargs for SESSION such that it can store TABLE_SIZE char units of
   target variable table, reusing cached kernargs allocations, if possible.  */

static inline struct kernargs *
allocate_session_kernargs (struct gomp_offload_session *session,
			   size_t table_size)
{
  GCN_DEBUG ("Session %p asked for allocation of kernargs+%llu...\n",
	     session, (unsigned long long) table_size);
  struct agent_info *agent = session->agent;
  assert (!session->kernarg_cache_node);

  /* To increase chance of cache hit, round up size of the target variable
     table to a multiple of (64*sizeof(void*)), and ensure that this size is
     nonzero.  */
  if (!table_size)
    table_size++;

  {
    constexpr size_t rounding_factor = 64 * sizeof (void*);
    table_size += rounding_factor - 1;
    table_size = (table_size / rounding_factor) * rounding_factor;
  }
  size_t kernargs_size = sizeof (struct kernargs) + table_size;

  session->kernarg_cache_node = (alloc_cache_try_find
				 (&agent->kernarg_cache,
				  kernargs_size));

  if (!session->kernarg_cache_node)
    {
      /* Cache miss.  */
      void *ka_addr;
      hsa_status_t status = hsa_fns.hsa_memory_allocate_fn
	(agent->kernarg_region, sizeof (struct kernargs), &ka_addr);
      if (status != HSA_STATUS_SUCCESS)
	hsa_fatal ("Could not allocate memory for GCN kernel arguments", status);

      session->kernarg_cache_node = (alloc_cache_add_taken_node
				     (&agent->kernarg_cache, ka_addr,
				      kernargs_size));
      if (!session->kernarg_cache_node)
	GOMP_PLUGIN_fatal ("Could not allocate cache node for kernel arguments");
    }

  return session->kernarg_cache_node->allocation;
}

void **
GOMP_OFFLOAD_session_allocate_target_var_table (struct gomp_offload_session *session,
						size_t table_size)
{
  GCN_DEBUG ("Session %p asked to allocate\n", session);
  /* libgomp wants us to handle the TVT.  */
  assert (!session->target_var_table);

  if (inhibit_kernargs_tvt_merge)
    /* ... but the user does not.  Used for testing.  */
    return NULL;

  struct kernargs *kernargs = allocate_session_kernargs (session, table_size);
  return session->target_var_table = &kernargs->target_variable_table[0];
}

void
GOMP_OFFLOAD_session_set_target_var_table (struct gomp_offload_session *session,
					   void **table)
{
  GCN_DEBUG ("Session %p will use TVT %p...\n", session, table);
  assert (!session->target_var_table);
  /* libgomp wants to handle the TVT.  */
  allocate_session_kernargs (session, 0);
  session->target_var_table = table;
}

/* Create or reuse a team arena and stack space.
 
   Team arenas are used by OpenMP to avoid calling malloc multiple times
   while setting up each team.  This is purely a performance optimization.

   The stack space is used by all kernels.  We must allocate it in such a
   way that the reverse offload implementation can access the data.

   Allocating this memory costs performance, so this function will reuse an
   existing allocation if a large enough one is idle.
   The memory lock is released, but not deallocated, when the kernel exits.  */

static void
configure_ephemeral_memories (struct kernel_info *kernel,
			      struct kernargs_abi *kernargs, int num_teams,
			      int num_threads)
{
  struct agent_info *agent = kernel->agent;
  struct ephemeral_memories_list **next_ptr = &agent->ephemeral_memories_list;
  struct ephemeral_memories_list *item;

  int actual_arena_size = (kernel->kind == KIND_OPENMP
			   ? team_arena_size : 0);
  int actual_arena_total_size = actual_arena_size * num_teams;
  size_t size = (actual_arena_total_size
		 + num_teams * num_threads * stack_size);

  for (item = *next_ptr; item; next_ptr = &item->next, item = item->next)
    {
      if (item->size < size)
	continue;

      if (pthread_mutex_trylock (&item->in_use) == 0)
	break;
    }

  if (!item)
    {
      GCN_DEBUG ("Creating a new %sstack for %d teams with %d threads"
		 " (%zd bytes)\n", (actual_arena_size ? "arena and " : ""),
		 num_teams, num_threads, size);

      if (pthread_mutex_lock (&agent->ephemeral_memories_write_lock))
	{
	  GOMP_PLUGIN_error ("Could not lock a GCN agent program mutex");
	  return;
	}
      item = malloc (sizeof (*item));
      item->size = size;
      item->next = NULL;
      *next_ptr = item;

      if (pthread_mutex_init (&item->in_use, NULL))
	{
	  GOMP_PLUGIN_error ("Failed to initialize a GCN memory write mutex");
	  return;
	}
      if (pthread_mutex_lock (&item->in_use))
	{
	  GOMP_PLUGIN_error ("Could not lock a GCN agent program mutex");
	  return;
	}
      if (pthread_mutex_unlock (&agent->ephemeral_memories_write_lock))
	{
	  GOMP_PLUGIN_error ("Could not unlock a GCN agent program mutex");
	  return;
	}

      hsa_status_t status;
      status = hsa_fns.hsa_memory_allocate_fn (agent->data_region, size,
					       &item->address);
      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->address, agent->id,
						   HSA_ACCESS_PERMISSION_RW);
      if (status != HSA_STATUS_SUCCESS)
	hsa_fatal ("Could not assign arena & stack memory to device", status);
    }

  kernargs->arena_ptr = (actual_arena_total_size
			 ? (uint64_t)item->address
			 : 0);
  kernargs->stack_ptr = (uint64_t)item->address + actual_arena_total_size;
  kernargs->arena_size_per_team = actual_arena_size;
  kernargs->stack_size_per_thread = stack_size;
}

/* Mark an ephemeral memory space available for reuse.  */

static void
release_ephemeral_memories (struct agent_info* agent, void *address)
{
  struct ephemeral_memories_list *item;

  for (item = agent->ephemeral_memories_list; item; item = item->next)
    {
      if (item->address == address)
	{
	  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_ephemeral_memories (struct agent_info *agent)
{
  struct ephemeral_memories_list *item, *next;

  for (item = agent->ephemeral_memories_list; item; item = next)
    {
      next = item->next;
      hsa_fns.hsa_memory_free_fn (item->address);
      if (pthread_mutex_destroy (&item->in_use))
	{
	  GOMP_PLUGIN_error ("Failed to destroy a GCN memory mutex");
	  return false;
	}
      free (item);
    }
  agent->ephemeral_memories_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);

  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 inline void
prepare_kernel_dispatch (struct kernel_dispatch *shadow,
			 struct kernel_info *kernel, int num_teams,
			 int num_threads, struct kernargs *kernargs)
{
  struct agent_info *agent = kernel->agent;

  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;
  shadow->private_segment_size = kernel->private_segment_size;

  if (lowlat_size < 0)
    {
      /* Divide the LDS between the number of running teams.
	 Allocate not less than is defined in the kernel metadata.  */
      int teams_per_cu = num_teams / get_cu_count (agent);
      int LDS_per_team = (teams_per_cu ? 65536 / teams_per_cu : 65536);
      shadow->group_segment_size
	= (kernel->group_segment_size > LDS_per_team
	   ? kernel->group_segment_size
	   : LDS_per_team);;
    }
  else if (lowlat_size < GCN_LOWLAT_HEAP+8)
    /* Ensure that there's space for the OpenMP libgomp data.  */
    shadow->group_segment_size = GCN_LOWLAT_HEAP+8;
  else
    shadow->group_segment_size = (lowlat_size > 65536
				  ? 65536
				  : lowlat_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;
    }

  /* Zero-initialize the output_data (minimum needed).  */
  kernargs->abi.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->abi.heap_ptr = (int64_t)kernel->module->heap;

  /* Create the ephemeral memory spaces.  */
  configure_ephemeral_memories (kernel, &kernargs->abi, num_teams, num_threads);

  /* Ensure we can recognize unset return values.  */
  kernargs->output_data.return_value = 0xcafe0000;
}

/* Copy information from DISPATCH into PACKET, to get it ready for
   dispatching.  */

static inline void
populate_packet_from_dispatch (hsa_kernel_dispatch_packet_t *packet,
			       struct kernel_dispatch *shadow)
{
  packet->private_segment_size = shadow->private_segment_size;
  packet->group_segment_size = shadow->group_segment_size;
  packet->kernel_object = shadow->object;
  packet->completion_signal = shadow->signal;
}

static void
process_reverse_offload (uint64_t fn, uint64_t mapnum, uint64_t hostaddrs,
			 uint64_t sizes, uint64_t kinds, uint64_t dev_num64)
{
  int dev_num = dev_num64;
  GOMP_PLUGIN_target_rev (fn, mapnum, hostaddrs, sizes, kinds, dev_num,
			  NULL);
}

/* 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;
	case 4:
	  process_reverse_offload (data->value_u64[0], data->value_u64[1],
				   data->value_u64[2], data->value_u64[3],
				   data->value_u64[4], data->value_u64[5]);
	  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 inline void
cleanup_kernel_dispatch (struct kernel_dispatch *shadow,
			 struct kernargs *kernargs)
{
  GCN_DEBUG ("Released kernel dispatch: %p\n", shadow);

  void *addr = (void *)kernargs->abi.arena_ptr;
  if (!addr)
    addr = (void *)kernargs->abi.stack_ptr;
  release_ephemeral_memories (shadow->agent, addr);

  hsa_fns.hsa_signal_destroy_fn (shadow->signal);
}

/* 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;
  char *buf = alloca (strlen (kernel->name) + 4);
  sprintf (buf, "%s.kd", kernel->name);
  status = hsa_fns.hsa_executable_get_symbol_fn (agent->executable, NULL,
						 buf, 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", buf);
      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 as part of SESSION 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 gomp_offload_session *session,
	    struct kernel_info *kernel,
	    struct GOMP_kernel_launch_attributes *kla,
	    struct goacc_asyncqueue *aq, bool module_locked)
{
  struct agent_info *agent = kernel->agent;
  GCN_DEBUG ("SGPRs: %d, VGPRs: %d\n", kernel->description->sgpr_count,
	     kernel->description->vpgr_count);

  /* Reduce the number of threads/workers if there are insufficient
     VGPRs available to run the kernels together.  */
  if (kla->ndim == 3 && kernel->description->vpgr_count > 0)
    {
      int max_vgprs = max_isa_vgprs (agent->device_isa);
      int granulated_vgprs = (kernel->description->vpgr_count + 3) & ~3;
      int max_threads = (max_vgprs / granulated_vgprs) * 4;
      if (kla->gdims[2] > max_threads)
	{
	  GCN_WARNING ("Too many VGPRs required to support %d threads/workers"
		       " per team/gang - reducing to %d threads/workers.\n",
		       kla->gdims[2], max_threads);
	  kla->gdims[2] = max_threads;
	}
    }

  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 ();

  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]);

  struct kernargs *kernargs = session->kernarg_cache_node->allocation;
  packet->kernarg_address = kernargs;

  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;
  prepare_kernel_dispatch (&shadow, kernel, packet->grid_size_x,
			   packet->grid_size_z, kernargs);
  shadow.queue = command_q;

  if (debug)
    {
      fprintf (stderr, "\nKernel has following dependencies:\n");
      print_kernel_dispatch (&shadow, 2, kernargs);
    }

  populate_packet_from_dispatch (packet, &shadow);

  hsa_signal_t s = shadow.signal;
  hsa_fns.hsa_signal_store_relaxed_fn (s, 1);

  GCN_DEBUG ("Copying kernel runtime pointer %p to kernarg_address\n", session->target_var_table);
  memcpy (kernargs, &session->target_var_table, sizeof (session->target_var_table));

  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, kernargs, false);
    }
  console_output (kernel, kernargs, true);

  unsigned int return_value = (unsigned int)kernargs->output_data.return_value;

  cleanup_kernel_dispatch (&shadow, kernargs);
  release_session (session);

  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;
  kernel->description = d;
  if (pthread_mutex_init (&kernel->init_mutex, NULL))
    {
      GOMP_PLUGIN_error ("Failed to initialize a GCN kernel mutex");
      return false;
    }
  return true;
}

/* If status is SUCCESS, assume that the code runs if either the ISA of agent
   and code is the same - or it is generic code.
   Otherwise, execution failed with the provided status code; try to give
   some useful diagnostic.  */

static bool
isa_matches_agent (struct agent_info *agent, Elf64_Ehdr *image,
		   hsa_status_t status)
{
  /* Generic image - assume that it works and only return to here
     when it fails, i.e. fatal == true.  */
  if (status == HSA_STATUS_SUCCESS && elf_gcn_isa_is_generic (image))
    return true;

  int isa_field = elf_gcn_isa_field (image);
  if (status == HSA_STATUS_SUCCESS && isa_field == agent->device_isa)
    return true;

  /* If we get here, either the binary is non-generic and has a mismatch of
     the ISA - or is generic but not handled by the ROCm (e.g. because ROCm
     is too old).  */

  char msg[340];
  char agent_isa_xs[8];
  char device_isa_xs[8];
  const char *agent_isa_s = isa_name (agent->device_isa);
  const char *device_isa_s = isa_name (isa_field);
  if (agent_isa_s == NULL)
    {
      snprintf (agent_isa_xs, sizeof agent_isa_xs,
		"0x%X", agent->device_isa);
      agent_isa_s = agent_isa_xs;
    }
  if (device_isa_s == NULL)
    {
      snprintf (device_isa_xs, sizeof device_isa_xs, "0x%X", isa_field);
      device_isa_s = device_isa_xs;
    }

  /* Some error which should be unrelated to the ISA.  */
  if (status != HSA_STATUS_SUCCESS
      && status != HSA_STATUS_ERROR_INVALID_CODE_OBJECT
      && status != HSA_STATUS_ERROR_INVALID_ISA_NAME
      && status != HSA_STATUS_ERROR_INCOMPATIBLE_ARGUMENTS)
    snprintf (msg, sizeof msg,
	      "Could not load GCN code object with ISA %s on GPU with "
	      "ISA %s (device %d).\n"
	      "Consider using ROCR_VISIBLE_DEVICES to disable incompatible "
	      "devices or run with LOADER_ENABLE_LOGGING=1 for more details.",
	      device_isa_s, agent_isa_s, agent->device_id);
  else if (status == HSA_STATUS_ERROR_INVALID_ISA_NAME
	   && elf_gcn_isa_is_generic (image))
    snprintf (msg, sizeof msg,
	      "Unsupported generic ISA %s on GPU with ISA %s (device %d).\n"
	      "%s%s%s"
	      "Consider using ROCR_VISIBLE_DEVICES to disable incompatible "
	      "devices, run with LOADER_ENABLE_LOGGING=1 for more details, "
	      "or try updating to a ROCm that supports this generic ISA.",
	      device_isa_s, agent_isa_s, agent->device_id,
	      agent_isa_s[0] != '0'
	      ? "Try to recompile with '-foffload-options=-march=" : "",
	      agent_isa_s[0] != '0' ? agent_isa_s : "",
	      agent_isa_s[0] != '0' ? ".\n" : "");
  else if (agent_isa_s[0] == '0')
    snprintf (msg, sizeof msg,
	      "GCN code object ISA '%s' is incompatible with GPU ISA '%s' "
	      "(device %d).\n"
	      "Consider using ROCR_VISIBLE_DEVICES to disable incompatible "
	      "devices or run with LOADER_ENABLE_LOGGING=1 for more details.",
	      device_isa_s, agent_isa_s, agent->device_id);
  else if (strcmp (device_isa_s, agent_isa_s) == 0
	   || (elf_gcn_isa_is_generic (image)
	       && generic_isa_code (agent->device_isa) == isa_field))
    snprintf (msg, sizeof msg,
	      "GCN code object features do not match for an unknown reason "
	      "(device %d).\n"
	      "Try to adjust the HSA_XNACK setting (perhaps?), or use\n"
	      "ROCR_VISIBLE_DEVICES to disable incompatible devices.\n",
	      agent->device_id);
  else
    snprintf (msg, sizeof msg,
	      "GCN code object ISA '%s' is incompatible with GPU ISA '%s' "
	      "(device %d).\n"
	      "Try to recompile with '-foffload-options=-march=%s',\n"
	      "or use ROCR_VISIBLE_DEVICES to disable incompatible "
	      "devices.\n",
	      device_isa_s, agent_isa_s, agent->device_id, agent_isa_s);

  hsa_error (msg, status != HSA_STATUS_SUCCESS
		  ? status : HSA_STATUS_ERROR_INVALID_CODE_OBJECT);
  return false;
}

/* 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;
  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;

      /* Check the ISA early because older ROCm had unhelpful errors.  */
      if (!isa_matches_agent (agent, image, HSA_STATUS_SUCCESS))
	goto fail;

      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)
	{
	  isa_matches_agent (agent, image, 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;
    }

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);
      struct gomp_offload_session session;
      init_session (&session, module->fini_array_func->agent);
      GOMP_OFFLOAD_session_set_target_var_table (&session, NULL);
      run_kernel (&session, module->fini_array_func, &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.session,
		  entry->u.launch.kernel,
		  &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.  The aq->mutex
   lock should be held on entry, and remains locked on exit.  */

static void
wait_for_queue_nonfull (struct goacc_asyncqueue *aq)
{
  if (aq->queue_n == ASYNC_QUEUE_SIZE)
    {
      /* 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);
    }
}

/* 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 gomp_offload_session *session,
		   struct kernel_info *kernel,
		   struct GOMP_kernel_launch_attributes *kla)
{
  assert (aq->agent == kernel->agent);

  pthread_mutex_lock (&aq->mutex);

  wait_for_queue_nonfull (aq);

  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.session = session;
  aq->queue[queue_last].u.launch.kernel = kernel;
  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)
{
  pthread_mutex_lock (&aq->mutex);

  wait_for_queue_nonfull (aq);

  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)
{
  pthread_mutex_lock (&aq->mutex);

  wait_for_queue_nonfull (aq);

  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;

  pthread_mutex_lock (&aq->mutex);

  wait_for_queue_nonfull (aq);

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

/* A wrapper that works around an issue in the HSA runtime with host-to-device
   copies from read-only pages.  */

static void
hsa_memory_copy_wrapper (void *dst, const void *src, size_t len)
{
  hsa_status_t status = hsa_fns.hsa_memory_copy_fn (dst, src, len);

  if (status == HSA_STATUS_SUCCESS)
    return;

  /* It appears that the copy fails if the source data is in a read-only page.
     We can't detect that easily, so try copying the data to a temporary buffer
     and doing the copy again if we got an error above.  */

  GCN_WARNING ("Read-only data transfer bug workaround triggered for "
	       "[%p:+%d]\n", (void *) src, (int) len);

  void *src_copy = malloc (len);
  memcpy (src_copy, src, len);
  status = hsa_fns.hsa_memory_copy_fn (dst, (const void *) src_copy, len);
  free (src_copy);
  if (status != HSA_STATUS_SUCCESS)
    GOMP_PLUGIN_error ("memory copy failed");
}

/* 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_memory_copy_wrapper (data->dst, data->src, data->len);
  free (data);
}

/* Request an asynchronous data copy, to or from a device, on a given queue.
   The event will be registered as a callback.  */

static void
queue_push_copy (struct goacc_asyncqueue *aq, void *dst, const void *src,
		 size_t len)
{
  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->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);
}

/* }}}  */
/* {{{ Managed Memory

   This implements an allocator equivalent to CUDA "Managed" memory, in which
   the pages automatically migrate between host and device memory, as needed.
   These allocations are visible from both the host and devices without the
   need for explicit mappings.  However, OpenMP does need "is_device_ptr" or
   "has_device_addr" to function properly.

   There isn't a high-level HSA/ROCr API to allocate managed memory, so we
   use regular memory and register it with the driver by setting it to
   "coarse-grained" mode, and setting the "accessible by default" attribute
   on devices where HSA_AMD_SYSTEM_INFO_SVM_ACCESSIBLE_BY_DEFAULT isn't set
   as standard (as it isn't on systems that don't support USM, or when
   HSA_XNACK != 1).

   This is in contrast to GOMP_OFFLOAD_alloc which allocates coarse-grained
   *GPU memory*, which is not visible on the host.

   It would be possible to register memory returned by malloc, but
   experimentation shows that doing so causes memory faults within the HSA
   runtime code.  Therefore, the Managed memory space is allocated as a
   largish block and then subdivided via a custom allocator.  The "simple"
   allocator is designed specifically to store its free-chain outside of
   the registered pages so that allocation does not inadvertently cause
   pages to migrate.

   Note: if the user has multiple mismatched devices, and one or more do
   not support USM (or XNACK is off), then each page of the Managed heap
   could end up associated with a different device (by calling omp_alloc
   before and after omp_set_default_device).  This issue remains
   an *unhandled* edge-case, at present.  */

gomp_simple_alloc_ctx_p managed_ctx = NULL;

/* Initialize or extend the Managed memory space.  This is called whenever
   allocation fails.  SIZE is the minimum size required for the failed
   allocation to succeed; the function may choose a larger size.
   Note that Linux lazy allocation means that the memory returned isn't
   guaranteed to actually exist.  */

static bool
managed_heap_create (struct agent_info *agent, size_t size)
{
  static int lock = 0;
  while (__atomic_exchange_n (&lock, 1, __ATOMIC_ACQUIRE) != 0)
    ;

  size_t default_size = 1L * 1024 * 1024 * 1024; /* 1GB */
  if (size < default_size)
    size = default_size;

  /* Round up to a whole page.  */
  int pagesize = getpagesize ();
  int misalignment = size % pagesize;
  if (misalignment > 0)
    size += pagesize - misalignment;

  /* Try to get contiguous memory, but it might not be possible.
     The most recent previous allocation is at the head of the list.  */
  static void *addrhint = NULL;
  void *new_pages = mmap (addrhint, size, PROT_READ | PROT_WRITE,
			  MAP_PRIVATE | MAP_ANONYMOUS, -1, 0);
  if (!new_pages)
    {
      GCN_DEBUG ("Could not allocate Managed Memory heap.");
      __atomic_store_n (&lock, 0, __ATOMIC_RELEASE);
      return false;
    }

  /* Register the heap allocation as coarse grained, "Managed" memory.  */
  struct hsa_amd_svm_attribute_pair_s attr = {
    HSA_AMD_SVM_ATTRIB_GLOBAL_FLAG,
    HSA_AMD_SVM_GLOBAL_FLAG_COARSE_GRAINED
  };
  hsa_status_t status = hsa_fns.hsa_amd_svm_attributes_set_fn (new_pages, size,
							       &attr, 1);
  if (status != HSA_STATUS_SUCCESS)
    GOMP_PLUGIN_fatal ("Failed to allocate Unified Shared Memory;"
		       " please update your drivers and/or kernel");

  /* The HSA_AMD_SVM_ATTRIB_AGENT_ACCESSIBLE setting is required on devices
     without default SVM.  */
  static int svm_accessible = 0xff; /* Use 0xff as "undefined".  */
  if (svm_accessible == 0xff)
    {
      status = hsa_fns.hsa_system_get_info_fn
	(HSA_AMD_SYSTEM_INFO_SVM_ACCESSIBLE_BY_DEFAULT, &svm_accessible);
      if (status != HSA_STATUS_SUCCESS)
	{
	  GCN_DEBUG ("warning: failed to query "
		     " HSA_AMD_SYSTEM_INFO_SVM_ACCESSIBLE_BY_DEFAULT\n");
	  svm_accessible = false;
	}
    }
  if (svm_accessible == false)
    {
      struct hsa_amd_svm_attribute_pair_s attr2;
      attr2.attribute = HSA_AMD_SVM_ATTRIB_AGENT_ACCESSIBLE;
      attr2.value = agent->id.handle;
      status = hsa_fns.hsa_amd_svm_attributes_set_fn (new_pages, size, &attr2,
						      1);
      if (status != HSA_STATUS_SUCCESS)
	GOMP_PLUGIN_fatal ("Failed to allocate Unified Shared Memory;"
			   " please update your drivers and/or kernel");
    }

  addrhint = new_pages + size;

  /* Initialize a new Managed memory heap, or add the new memory into an
     existing Managed memory heap.  */
  if (!managed_ctx)
    managed_ctx = gomp_simple_alloc_init_context ();
  gomp_simple_alloc_register_memory (managed_ctx, new_pages, size);

  __atomic_store_n (&lock, 0, __ATOMIC_RELEASE);
  return true;
}

/* }}} */
/* {{{ OpenACC support  */

/* Execute an OpenACC kernel, synchronously or asynchronously.  */

static void
gcn_exec (struct kernel_info *kernel, struct gomp_offload_session *session,
	  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;

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

  /* Ideally, when a dimension isn't explicitly specified, we should
     tune it to run 40 (or 32?) threads per CU with no threads getting queued.
     In practice, we tune for peak performance on BabelStream, which
     for OpenACC is currently 32 threads per CU.  */
  if (dims[0] == 0 && dims[1] == 0)
    {
      /* 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 workers
	 and gangs.  */

      dims[0] = get_cu_count (kernel->agent) * 4; /* Gangs.  */
      dims[1] = 8; /* Workers.  */
    }
  else if (dims[0] == 0 && dims[1] > 0)
    {
      /* Auto-scale the number of gangs with the requested number of workers.  */
      dims[0] = get_cu_count (kernel->agent) * (32 / dims[1]);
    }
  else if (dims[0] > 0 && dims[1] == 0)
    {
      /* Auto-scale the number of workers with the requested number of gangs.  */
      dims[1] = get_cu_count (kernel->agent) * 32 / dims[0];
      if (dims[1] == 0)
	dims[1] = 1;
      if (dims[1] > 16)
	dims[1] = 16;
    }

  /* 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 (session, kernel, &kla, NULL, false);
  else
    queue_push_launch (aq, session, kernel, &kla);

  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  */

#if 0  /* TODO: Use to enable self-mapping/USM automatically.  */
/* FIXME: The auto-self-map feature depends on still mapping 'declare target'
   variables, even if ignoring all other mappings. Cf. PR 115279.  */

/* Return TRUE if the GPU is an APU, i.e. the GPU is integrated with the CPU
   such that both use the same memory controller such that mapping or memory
   migration is pointless.  If CHECK_XNACK is TRUE, it additionally requires
   that the GPU has *no* XNACK support otherwise FALSE is returned.

   In theory, enabling unified-shared memory for APUs should always work,
   however, with AMD GPUs some APUs (e.g. MI300A) still require XNACK to be
   enabled as it is required to handle page faults.

   Thus, for unified-shared memory access, either of the following must hold:
   * HSA_AMD_SYSTEM_INFO_SVM_ACCESSIBLE_BY_DEFAULT is TRUE
     This implies that all GPUs support USM access, either directly (as APU)
     or via page migration.  For MI300A, this is only the case if
     HSA_AMD_SYSTEM_INFO_XNACK_ENABLED is TRUE.
   * If the GPU an APU *and* it does not support XNACK.  */

static bool
is_integrated_apu (struct agent_info *agent, bool check_xnack)
{
  enum {
    HSACO_ATTR_UNSUPPORTED,
    HSACO_ATTR_OFF,
    HSACO_ATTR_ON,
    HSACO_ATTR_ANY,
    HSACO_ATTR_DEFAULT
  };

  bool is_apu;
  uint8_t mem_prop[8];
  hsa_status_t status;

  status = hsa_fns.hsa_agent_get_info_fn (
	     agent->id, (hsa_agent_info_t) HSA_AMD_AGENT_INFO_MEMORY_PROPERTIES,
	     mem_prop);
  _Static_assert (HSA_AMD_MEMORY_PROPERTY_AGENT_IS_APU < 8,
		  "HSA_AMD_MEMORY_PROPERTY_AGENT_IS_APU < 8");
  is_apu = (status == HSA_STATUS_SUCCESS
	    && (mem_prop[0] & (1 << HSA_AMD_MEMORY_PROPERTY_AGENT_IS_APU)));

  if (check_xnack)
    switch(agent->device_isa)
      {
#define GCN_DEVICE(name, NAME, ELF, ISA, XNACK, ...) \
      case ELF: return is_apu && (XNACK == HSACO_ATTR_UNSUPPORTED);
#include "../../gcc/config/gcn/gcn-devices.def"
      default: return false;  /* Just to be save.  */
      }
  return is_apu;
}
#endif

/* Return the name of the accelerator, which is "gcn".  */

const char *
GOMP_OFFLOAD_get_name (void)
{
  return "gcn";
}

/* Return the UID; if not available return NULL.
   Returns freshly allocated memory.  */

const char *
GOMP_OFFLOAD_get_uid (int ord)
{
  char *str;
  hsa_status_t status;
  struct agent_info *agent = get_agent_info (ord);

  /* HSA documentation states: maximally 21 characters including NUL.  */
  str = GOMP_PLUGIN_malloc (21 * sizeof (char));
  status = hsa_fns.hsa_agent_get_info_fn (agent->id, HSA_AMD_AGENT_INFO_UUID,
					  str);
  if (status != HSA_STATUS_SUCCESS)
    {
      free (str);
      return NULL;
    }
  return str;
}

/* 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 (unsigned int omp_requires_mask)
{
  if (!init_hsa_context (true))
    exit (EXIT_FAILURE);
  /* Return -1 if no omp_requires_mask cannot be fulfilled but
     devices were present.  */
  if (hsa_context.agent_count > 0
      && ((omp_requires_mask
	   & ~(GOMP_REQUIRES_UNIFIED_ADDRESS
	       | GOMP_REQUIRES_UNIFIED_SHARED_MEMORY
	       | GOMP_REQUIRES_SELF_MAPS
	       | GOMP_REQUIRES_REVERSE_OFFLOAD)) != 0))
    return -1;
  /* Check whether host page access is supported; this is per system level
     (all GPUs supported by HSA).  While intrinsically true for APUs, it
     requires XNACK support for discrete GPUs.  */
  if (hsa_context.agent_count > 0
      && (omp_requires_mask
	  & (GOMP_REQUIRES_UNIFIED_SHARED_MEMORY | GOMP_REQUIRES_SELF_MAPS)))
    {
      bool b;
      hsa_system_info_t type = HSA_AMD_SYSTEM_INFO_SVM_ACCESSIBLE_BY_DEFAULT;
      hsa_status_t status = hsa_fns.hsa_system_get_info_fn (type, &b);
      if (status != HSA_STATUS_SUCCESS)
	GOMP_PLUGIN_error ("HSA_AMD_SYSTEM_INFO_SVM_ACCESSIBLE_BY_DEFAULT "
			   "failed");
      if (!b)
	return -1;
    }

  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 (false))
    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->ephemeral_memories_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->ephemeral_memories_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);

  status = hsa_fns.hsa_agent_get_info_fn (agent->id, HSA_AGENT_INFO_NAME,
					  &agent->name);
  if (status != HSA_STATUS_SUCCESS)
    return hsa_error ("Error querying the name of the agent", status);

  agent->device_isa = isa_code (agent->name);
  if (agent->device_isa == EF_AMDGPU_MACH_UNSUPPORTED)
    {
      char msg[33 + 64 + 1];
      snprintf (msg, sizeof msg,
		"Unknown GCN agent architecture '%s'", agent->name);
      return hsa_error (msg, HSA_STATUS_ERROR);
    }

  status = hsa_fns.hsa_agent_get_info_fn (agent->id, HSA_AGENT_INFO_VENDOR_NAME,
					  &agent->vendor_name);
  if (status != HSA_STATUS_SUCCESS)
    return hsa_error ("Error querying the vendor name of the agent", status);

  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 (status != HSA_STATUS_SUCCESS
      && status != HSA_STATUS_INFO_BREAK)
    hsa_error ("Scanning memory regions failed", status);
  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 (status != HSA_STATUS_SUCCESS
      && status != HSA_STATUS_INFO_BREAK)
    hsa_error ("Scanning memory regions failed", status);
  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);

  /* Prepare kernargs cache.  */
  init_alloc_cache (&agent->kernarg_cache);

  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.  If not NULL, REV_FN_TABLE will
   contain the on-device addresses of the functions for reverse offload.  To be
   freed by the caller.  */

int
GOMP_OFFLOAD_load_image (int ord, unsigned version, const void *target_data,
			 struct addr_pair **target_table,
			 uint64_t **rev_fn_table,
			 uint64_t *host_ind_fn_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 ind_func_count = GOMP_VERSION_SUPPORTS_INDIRECT_FUNCS (version)
			      ? image_desc->ind_func_count : 0;
  unsigned var_count = image_desc->global_variable_count;
  /* Currently, "others" is a struct of ICVS.  */
  int other_count = 1;

  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 %d indirect functions in an image\n", ind_func_count);
  GCN_DEBUG ("Encountered %u global variables in an image\n", var_count);
  GCN_DEBUG ("Expect %d other variables in an image\n", other_count);
  pair = GOMP_PLUGIN_malloc ((kernel_count + var_count + other_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;

  if (var_count > 0)
    {
      hsa_status_t status;
      hsa_executable_symbol_t var_symbol;
      status = hsa_fns.hsa_executable_get_symbol_fn (agent->executable, NULL,
						     ".offload_var_table",
						     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_table_addr;
      status = hsa_fns.hsa_executable_symbol_get_info_fn
	(var_symbol, HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_ADDRESS,
	 &var_table_addr);
      if (status != HSA_STATUS_SUCCESS)
	hsa_fatal ("Could not extract a variable from its symbol", status);

      struct {
	uint64_t addr;
	uint64_t size;
      } var_table[var_count];
      GOMP_OFFLOAD_dev2host (agent->device_id, var_table,
			     (void*)var_table_addr, sizeof (var_table));

      for (unsigned i = 0; i < var_count; i++)
	{
	  pair->start = var_table[i].addr;
	  pair->end = var_table[i].addr + var_table[i].size;
	  GCN_DEBUG ("Found variable at %p with size %lu\n",
		     (void *)var_table[i].addr, var_table[i].size);
	  pair++;
	}
    }

  if (ind_func_count > 0)
    {
      hsa_status_t status;

      /* Read indirect function table from image.  */
      hsa_executable_symbol_t ind_funcs_symbol;
      status = hsa_fns.hsa_executable_get_symbol_fn (agent->executable, NULL,
						     ".offload_ind_func_table",
						     agent->id,
						     0, &ind_funcs_symbol);

      if (status != HSA_STATUS_SUCCESS)
	hsa_fatal ("Could not find .offload_ind_func_table symbol in the "
		   "code object", status);

      uint64_t ind_funcs_table_addr;
      status = hsa_fns.hsa_executable_symbol_get_info_fn
	(ind_funcs_symbol, HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_ADDRESS,
	 &ind_funcs_table_addr);
      if (status != HSA_STATUS_SUCCESS)
	hsa_fatal ("Could not extract a variable from its symbol", status);

      uint64_t ind_funcs_table[ind_func_count];
      GOMP_OFFLOAD_dev2host (agent->device_id, ind_funcs_table,
			     (void*) ind_funcs_table_addr,
			     sizeof (ind_funcs_table));

      /* For newer binaries, the hash table for 'indirect' is created on the
	 host. Older binaries don't have GOMP_INDIRECT_ADDR_HMAP on the
	 device side - and have to create the table themselves using
	 GOMP_INDIRECT_ADDR_MAP.  */

      hsa_executable_symbol_t symbol;
      bool host_init_htab = true;
      #ifdef USE_HASHTAB_LOOKUP_FOR_INDIRECT
      status
	= hsa_fns.hsa_executable_get_symbol_fn (agent->executable, NULL,
						XSTRING (GOMP_INDIRECT_ADDR_HMAP),
						agent->id, 0, &symbol);
      if (status != HSA_STATUS_SUCCESS)
      #endif
	{
	  host_init_htab = false;
	  status = hsa_fns.hsa_executable_get_symbol_fn (agent->executable, NULL,
		     XSTRING (GOMP_INDIRECT_ADDR_MAP), agent->id, 0, &symbol);
	}
      if (status != HSA_STATUS_SUCCESS)
	hsa_fatal ("Could not find GOMP_INDIRECT_ADDR_MAP in code object",
		   status);
      uint64_t varptr;
      uint32_t varsize;

      status = hsa_fns.hsa_executable_symbol_get_info_fn
	(symbol, HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_ADDRESS,
	 &varptr);
      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
	(symbol, HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_SIZE,
	&varsize);
      if (status != HSA_STATUS_SUCCESS)
	hsa_fatal ("Could not extract a variable size from its symbol",
		   status);

      GCN_DEBUG ("Found GOMP_INDIRECT_ADDR_%sMAP at %lx with size %d\n",
		 host_init_htab ? "H" : "", varptr, varsize);

      void *map_target_addr;
      if (!host_init_htab)
	{
	  /* Build host->target address map for indirect functions.  */
	  uint64_t ind_fn_map[ind_func_count * 2 + 1];
	  for (unsigned i = 0; i < ind_func_count; i++)
	    {
	      ind_fn_map[i * 2] = host_ind_fn_table[i];
	      ind_fn_map[i * 2 + 1] = ind_funcs_table[i];
	      GCN_DEBUG ("Indirect function %d: %lx->%lx\n",
			 i, host_ind_fn_table[i], ind_funcs_table[i]);
	    }
	  ind_fn_map[ind_func_count * 2] = 0;
	  /* Write the map onto the target.  */
	  map_target_addr = GOMP_OFFLOAD_alloc (agent->device_id,
						sizeof (ind_fn_map));
	  GOMP_OFFLOAD_host2dev (agent->device_id, map_target_addr,
				 (void*) ind_fn_map, sizeof (ind_fn_map));
	}
      #ifdef USE_HASHTAB_LOOKUP_FOR_INDIRECT
      else
	{
	  /* FIXME: Handle multi-kernel load and unload, cf. PR 114690.  */
	  size_t host_map_size;
	  void *host_map;
	  host_map = create_target_indirect_map (&host_map_size, ind_func_count,
						 host_ind_fn_table,
						 ind_funcs_table);
	  for (unsigned i = 0; i < ind_func_count; i++)
	      GCN_DEBUG ("Indirect function %d: %lx->%lx\n",
			 i, host_ind_fn_table[i], ind_funcs_table[i]);
	  /* Write the map onto the target.  */
	  map_target_addr = GOMP_OFFLOAD_alloc (agent->device_id,
						host_map_size);
	  GOMP_OFFLOAD_host2dev (agent->device_id, map_target_addr,
				 host_map, host_map_size);
	}
      #endif

      GCN_DEBUG ("Allocated indirect map at %p\n", map_target_addr);

      /* Write address of the map onto the target.  */
      GOMP_OFFLOAD_host2dev (agent->device_id, (void *) varptr,
			     &map_target_addr,
			     sizeof (map_target_addr));
    }

  GCN_DEBUG ("Looking for variable %s\n", XSTRING (GOMP_ADDITIONAL_ICVS));

  hsa_status_t status;
  hsa_executable_symbol_t var_symbol;
  status = hsa_fns.hsa_executable_get_symbol_fn (agent->executable, NULL,
						 XSTRING (GOMP_ADDITIONAL_ICVS),
						 agent->id, 0, &var_symbol);
  if (status == HSA_STATUS_SUCCESS)
    {
      uint64_t varptr;
      uint32_t varsize;

      status = hsa_fns.hsa_executable_symbol_get_info_fn
	(var_symbol, HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_ADDRESS,
	 &varptr);
      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,
	 &varsize);
      if (status != HSA_STATUS_SUCCESS)
	hsa_fatal ("Could not extract a variable size from its symbol",
		   status);

      pair->start = varptr;
      pair->end = varptr + varsize;
    }
  else
    {
      /* The variable was not in this image.  */
      GCN_DEBUG ("Variable not found in image: %s\n",
		 XSTRING (GOMP_ADDITIONAL_ICVS));
      pair->start = pair->end = 0;
    }

  /* 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);
      struct gomp_offload_session session;
      init_session (&session, agent);
      GOMP_OFFLOAD_session_set_target_var_table (&session, NULL);
      run_kernel (&session, module->init_array_func, &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--;

  if (rev_fn_table != NULL && kernel_count == 0)
    *rev_fn_table = NULL;
  else if (rev_fn_table != NULL)
    {
      hsa_status_t status;
      hsa_executable_symbol_t var_symbol;
      status = hsa_fns.hsa_executable_get_symbol_fn (agent->executable, NULL,
						     ".offload_func_table",
						     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 fn_table_addr;
      status = hsa_fns.hsa_executable_symbol_get_info_fn
	(var_symbol, HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_ADDRESS,
	 &fn_table_addr);
      if (status != HSA_STATUS_SUCCESS)
	hsa_fatal ("Could not extract a variable from its symbol", status);
      *rev_fn_table = GOMP_PLUGIN_malloc (kernel_count * sizeof (uint64_t));
      GOMP_OFFLOAD_dev2host (agent->device_id, *rev_fn_table,
			     (void*) fn_table_addr,
			     kernel_count * sizeof (uint64_t));
    }

  return kernel_count + var_count + other_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_ephemeral_memories (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);

  /* Clean up kernargs cache.  */
  struct alloc_cache_node *node = agent->kernarg_cache.head;
  while (node)
    {
      hsa_fns.hsa_memory_free_fn (node->allocation);

      struct alloc_cache_node *curr_node = node;
      node = curr_node->next;
      destroy_alloc_cache_node (curr_node);
    }

  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->ephemeral_memories_write_lock))
    {
      GOMP_PLUGIN_error ("Failed to destroy a GCN memory 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)
    GOMP_PLUGIN_fatal ("kernel initialization failed");

  return true;
}

/* 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_status_t status = hsa_fns.hsa_memory_copy_fn (dst, src, n);
  if (status != HSA_STATUS_SUCCESS)
    GOMP_PLUGIN_error ("memory copy failed");
  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_memory_copy_wrapper (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);
      if (!agent->omp_async_queue)
	return false;
      queue_push_copy (agent->omp_async_queue, dst, src, n);
      return true;
    }

  GCN_DEBUG ("Copying %zu bytes from device %d (%p) to device %d (%p)\n", n,
	     device, src, device, dst);
  hsa_status_t status = hsa_fns.hsa_memory_copy_fn (dst, src, n);
  if (status != HSA_STATUS_SUCCESS)
    GOMP_PLUGIN_error ("memory copy failed");
  return true;
}

/* Here <quantity>_size refers to <quantity> multiplied by size -- i.e.
   measured in bytes.  So we have:

   dim1_size: number of bytes to copy on innermost dimension ("row")
   dim0_len: number of rows to copy
   dst: base pointer for destination of copy
   dst_offset1_size: innermost row offset (for dest), in bytes
   dst_offset0_len: offset, number of rows (for dest)
   dst_dim1_size: whole-array dest row length, in bytes (pitch)
   src: base pointer for source of copy
   src_offset1_size: innermost row offset (for source), in bytes
   src_offset0_len: offset, number of rows (for source)
   src_dim1_size: whole-array source row length, in bytes (pitch)
*/

int
GOMP_OFFLOAD_memcpy2d (int dst_ord, int src_ord, size_t dim1_size,
		       size_t dim0_len, void *dst, size_t dst_offset1_size,
		       size_t dst_offset0_len, size_t dst_dim1_size,
		       const void *src, size_t src_offset1_size,
		       size_t src_offset0_len, size_t src_dim1_size)
{
  if (!hsa_fns.hsa_amd_memory_lock_fn
      || !hsa_fns.hsa_amd_memory_unlock_fn
      || !hsa_fns.hsa_amd_memory_async_copy_rect_fn)
    return -1;

  /* GCN hardware requires 4-byte alignment for base addresses & pitches.  Bail
     out quietly if we have anything oddly-aligned rather than letting the
     driver raise an error.  */
  if ((((uintptr_t) dst) & 3) != 0 || (((uintptr_t) src) & 3) != 0)
    return -1;

  if ((dst_dim1_size & 3) != 0 || (src_dim1_size & 3) != 0)
    return -1;

  /* Only handle host to device or device to host transfers here.  */
  if ((dst_ord == -1 && src_ord == -1)
      || (dst_ord != -1 && src_ord != -1))
    return -1;

  hsa_amd_copy_direction_t dir
    = (src_ord == -1) ? hsaHostToDevice : hsaDeviceToHost;
  hsa_agent_t copy_agent;

  /* We need to pin (lock) host memory before we start the transfer.  Try to
     lock the minimum size necessary, i.e. using partial first/last rows of the
     whole array.  Something like this:

	 rows -->
	 ..............
     c | ..#######+++++ <- first row apart from {src,dst}_offset1_size
     o | ++#######+++++ <- whole row
     l | ++#######+++++ <-     "
     s v ++#######..... <- last row apart from trailing remainder
	 ..............

     We could split very large transfers into several rectangular copies, but
     that is unimplemented for now.  */

  size_t bounded_size_host, first_elem_offset_host;
  void *host_ptr;
  if (dir == hsaHostToDevice)
    {
      bounded_size_host = src_dim1_size * (dim0_len - 1) + dim1_size;
      first_elem_offset_host = src_offset0_len * src_dim1_size
			       + src_offset1_size;
      host_ptr = (void *) src;
      struct agent_info *agent = get_agent_info (dst_ord);
      copy_agent = agent->id;
    }
  else
    {
      bounded_size_host = dst_dim1_size * (dim0_len - 1) + dim1_size;
      first_elem_offset_host = dst_offset0_len * dst_dim1_size
			       + dst_offset1_size;
      host_ptr = dst;
      struct agent_info *agent = get_agent_info (src_ord);
      copy_agent = agent->id;
    }

  void *agent_ptr;

  hsa_status_t status
    = hsa_fns.hsa_amd_memory_lock_fn (host_ptr + first_elem_offset_host,
				      bounded_size_host, NULL, 0, &agent_ptr);
  /* We can't lock the host memory: don't give up though, we might still be
     able to use the slow path in our caller.  So, don't make this an
     error.  */
  if (status != HSA_STATUS_SUCCESS)
    return -1;

  hsa_pitched_ptr_t dstpp, srcpp;
  hsa_dim3_t dst_offsets, src_offsets, ranges;

  int retval = 1;

  hsa_signal_t completion_signal;
  status = hsa_fns.hsa_signal_create_fn (1, 0, NULL, &completion_signal);
  if (status != HSA_STATUS_SUCCESS)
    {
      retval = -1;
      goto unlock;
    }

  if (dir == hsaHostToDevice)
    {
      srcpp.base = agent_ptr - first_elem_offset_host;
      dstpp.base = dst;
    }
  else
    {
      srcpp.base = (void *) src;
      dstpp.base = agent_ptr - first_elem_offset_host;
    }

  srcpp.pitch = src_dim1_size;
  srcpp.slice = 0;

  src_offsets.x = src_offset1_size;
  src_offsets.y = src_offset0_len;
  src_offsets.z = 0;

  dstpp.pitch = dst_dim1_size;
  dstpp.slice = 0;

  dst_offsets.x = dst_offset1_size;
  dst_offsets.y = dst_offset0_len;
  dst_offsets.z = 0;

  ranges.x = dim1_size;
  ranges.y = dim0_len;
  ranges.z = 1;

  status
    = hsa_fns.hsa_amd_memory_async_copy_rect_fn (&dstpp, &dst_offsets, &srcpp,
						 &src_offsets, &ranges,
						 copy_agent, dir, 0, NULL,
						 completion_signal);
  /* If the rectangular copy fails, we might still be able to use the slow
     path.  We need to unlock the host memory though, so don't return
     immediately.  */
  if (status != HSA_STATUS_SUCCESS)
    retval = -1;
  else
    hsa_fns.hsa_signal_wait_acquire_fn (completion_signal,
					HSA_SIGNAL_CONDITION_LT, 1, UINT64_MAX,
					HSA_WAIT_STATE_ACTIVE);

  hsa_fns.hsa_signal_destroy_fn (completion_signal);

unlock:
  status = hsa_fns.hsa_amd_memory_unlock_fn (host_ptr + first_elem_offset_host);
  if (status != HSA_STATUS_SUCCESS)
    hsa_fatal ("Could not unlock host memory", status);

  return retval;
}

/* As above, <quantity>_size refers to <quantity> multiplied by size -- i.e.
   measured in bytes.  So we have:

   dim2_size: number of bytes to copy on innermost dimension ("row")
   dim1_len: number of rows per slice to copy
   dim0_len: number of slices to copy
   dst: base pointer for destination of copy
   dst_offset2_size: innermost row offset (for dest), in bytes
   dst_offset1_len: offset, number of rows (for dest)
   dst_offset0_len: offset, number of slices (for dest)
   dst_dim2_size: whole-array dest row length, in bytes (pitch)
   dst_dim1_len: whole-array number of rows in slice (for dest)
   src: base pointer for source of copy
   src_offset2_size: innermost row offset (for source), in bytes
   src_offset1_len: offset, number of rows (for source)
   src_offset0_len: offset, number of slices (for source)
   src_dim2_size: whole-array source row length, in bytes (pitch)
   src_dim1_len: whole-array number of rows in slice (for source)
*/

int
GOMP_OFFLOAD_memcpy3d (int dst_ord, int src_ord, size_t dim2_size,
		       size_t dim1_len, size_t dim0_len, void *dst,
		       size_t dst_offset2_size, size_t dst_offset1_len,
		       size_t dst_offset0_len, size_t dst_dim2_size,
		       size_t dst_dim1_len, const void *src,
		       size_t src_offset2_size, size_t src_offset1_len,
		       size_t src_offset0_len, size_t src_dim2_size,
		       size_t src_dim1_len)
{
  if (!hsa_fns.hsa_amd_memory_lock_fn
      || !hsa_fns.hsa_amd_memory_unlock_fn
      || !hsa_fns.hsa_amd_memory_async_copy_rect_fn)
    return -1;

  /* GCN hardware requires 4-byte alignment for base addresses & pitches.  Bail
     out quietly if we have anything oddly-aligned rather than letting the
     driver raise an error.  */
  if ((((uintptr_t) dst) & 3) != 0 || (((uintptr_t) src) & 3) != 0)
    return -1;

  if ((dst_dim2_size & 3) != 0 || (src_dim2_size & 3) != 0)
    return -1;

  /* Only handle host to device or device to host transfers here.  */
  if ((dst_ord == -1 && src_ord == -1)
      || (dst_ord != -1 && src_ord != -1))
    return -1;

  hsa_amd_copy_direction_t dir
    = (src_ord == -1) ? hsaHostToDevice : hsaDeviceToHost;
  hsa_agent_t copy_agent;

  /* We need to pin (lock) host memory before we start the transfer.  Try to
     lock the minimum size necessary, i.e. using partial first/last slices of
     the whole 3D array.  Something like this:

	 slice 0:          slice 1:	     slice 2:
	     __________        __________        __________
       ^    /+++++++++/    :  /+++++++++/    :	/         /
    column /+++##++++/|    | /+++##++++/|    | /+++##    /  # = subarray
     /	  /   ##++++/ |    |/+++##++++/ |    |/+++##++++/   + = area to pin
	 /_________/  :    /_________/  :    /_________/
	row --->

     We could split very large transfers into several rectangular copies, but
     that is unimplemented for now.  */

  size_t bounded_size_host, first_elem_offset_host;
  void *host_ptr;
  if (dir == hsaHostToDevice)
    {
      size_t slice_bytes = src_dim2_size * src_dim1_len;
      bounded_size_host = slice_bytes * (dim0_len - 1)
			  + src_dim2_size * (dim1_len - 1)
			  + dim2_size;
      first_elem_offset_host = src_offset0_len * slice_bytes
			       + src_offset1_len * src_dim2_size
			       + src_offset2_size;
      host_ptr = (void *) src;
      struct agent_info *agent = get_agent_info (dst_ord);
      copy_agent = agent->id;
    }
  else
    {
      size_t slice_bytes = dst_dim2_size * dst_dim1_len;
      bounded_size_host = slice_bytes * (dim0_len - 1)
			  + dst_dim2_size * (dim1_len - 1)
			  + dim2_size;
      first_elem_offset_host = dst_offset0_len * slice_bytes
			       + dst_offset1_len * dst_dim2_size
			       + dst_offset2_size;
      host_ptr = dst;
      struct agent_info *agent = get_agent_info (src_ord);
      copy_agent = agent->id;
    }

  void *agent_ptr;

  hsa_status_t status
    = hsa_fns.hsa_amd_memory_lock_fn (host_ptr + first_elem_offset_host,
				      bounded_size_host, NULL, 0, &agent_ptr);
  /* We can't lock the host memory: don't give up though, we might still be
     able to use the slow path in our caller (maybe even with iterated memcpy2d
     calls).  So, don't make this an error.  */
  if (status != HSA_STATUS_SUCCESS)
    return -1;

  hsa_pitched_ptr_t dstpp, srcpp;
  hsa_dim3_t dst_offsets, src_offsets, ranges;

  int retval = 1;

  hsa_signal_t completion_signal;
  status = hsa_fns.hsa_signal_create_fn (1, 0, NULL, &completion_signal);
  if (status != HSA_STATUS_SUCCESS)
    {
      retval = -1;
      goto unlock;
    }

  if (dir == hsaHostToDevice)
    {
      srcpp.base = agent_ptr - first_elem_offset_host;
      dstpp.base = dst;
    }
  else
    {
      srcpp.base = (void *) src;
      dstpp.base = agent_ptr - first_elem_offset_host;
    }

  /* Pitch is measured in bytes.  */
  srcpp.pitch = src_dim2_size;
  /* Slice is also measured in bytes (i.e. total per-slice).  */
  srcpp.slice = src_dim2_size * src_dim1_len;

  src_offsets.x = src_offset2_size;
  src_offsets.y = src_offset1_len;
  src_offsets.z = src_offset0_len;

  /* As above.  */
  dstpp.pitch = dst_dim2_size;
  dstpp.slice = dst_dim2_size * dst_dim1_len;

  dst_offsets.x = dst_offset2_size;
  dst_offsets.y = dst_offset1_len;
  dst_offsets.z = dst_offset0_len;

  ranges.x = dim2_size;
  ranges.y = dim1_len;
  ranges.z = dim0_len;

  status
    = hsa_fns.hsa_amd_memory_async_copy_rect_fn (&dstpp, &dst_offsets, &srcpp,
						 &src_offsets, &ranges,
						 copy_agent, dir, 0, NULL,
						 completion_signal);
  /* If the rectangular copy fails, we might still be able to use the slow
     path.  We need to unlock the host memory though, so don't return
     immediately.  */
  if (status != HSA_STATUS_SUCCESS)
    retval = -1;
  else
    {
      hsa_signal_value_t sv
	= hsa_fns.hsa_signal_wait_acquire_fn (completion_signal,
					      HSA_SIGNAL_CONDITION_LT, 1,
					      UINT64_MAX,
					      HSA_WAIT_STATE_ACTIVE);
      if (sv < 0)
	{
	  GCN_WARNING ("async copy rect failure");
	  retval = -1;
	}
    }

  hsa_fns.hsa_signal_destroy_fn (completion_signal);

unlock:
  status = hsa_fns.hsa_amd_memory_unlock_fn (host_ptr + first_elem_offset_host);
  if (status != HSA_STATUS_SUCCESS)
    hsa_fatal ("Could not unlock host memory", status);

  return retval;
}


static bool
init_hip_runtime_functions (void)
{
  bool inited = false;
  if (inited)
    return hip_fns.hipStreamCreate_fn != NULL;
  inited = true;

  void *handle = dlopen (hip_runtime_lib, RTLD_LAZY);
  if (handle == NULL)
    return false;

#define DLSYM_OPT_FN(function) \
  hip_fns.function##_fn = dlsym (handle, #function)

  DLSYM_OPT_FN (hipStreamCreate);
  DLSYM_OPT_FN (hipStreamDestroy);
  DLSYM_OPT_FN (hipStreamSynchronize);
  DLSYM_OPT_FN (hipCtxGetCurrent);
  DLSYM_OPT_FN (hipGetDevice);
  DLSYM_OPT_FN (hipSetDevice);
#undef DLSYM_OPT_FN

  if (!hip_fns.hipStreamCreate_fn
      || !hip_fns.hipStreamDestroy_fn
      || !hip_fns.hipStreamSynchronize_fn
      || !hip_fns.hipCtxGetCurrent_fn
      || !hip_fns.hipGetDevice_fn
      || !hip_fns.hipSetDevice_fn)
    {
      hip_fns.hipStreamCreate_fn = NULL;
      return false;
    }

  return true;
}

bool
GOMP_OFFLOAD_memset (int ord, void *ptr, int val, size_t count)
{
  hsa_status_t status = HSA_STATUS_SUCCESS;

  /* A memset feature is only provided via hsa_amd_memory_fill; while it
     is fast, it is an HSA extension and it has two requirements: The memory
     must be aligned to multiples of 4 bytes - and, by construction, only
     multiples of 4 bytes can be filled (uint32_t value argument).

     This means: Either not using that function or up to three function calls:
     - copy 1 to 3 bytes to get alignment (hsa_memory_copy), if unaligned
     - call hsa_amd_memory_fill
     - copy remaining 1 to 3 bytes (hsa_memory_copy), if after alignment
       count is not a multiple of 4 bytes.

     Having more than one function call is only profitable if there is
     enough data to process; see below for the used heuristic values.  */

  uint8_t v8 = (uint8_t) val;
  size_t before = (4 - (uintptr_t) ptr % 4) % 4;  /* 0 to 3 bytes.  */
  size_t tail = (count - before) % 4;  /* 0 to 3 bytes.  */

  /* Heuristic  */
  enum {
    /* Prefer alloca to malloc up to ... */
    alloca_size = 256,  /* bytes */
    /* Call hsa_amd_memory_fill also when two copy calls are required.  */
    always_use_fill = 256*1024,  /* bytes */
    /* Call hsa_amd_memory_fill also when on copy call is required.  */
    use_fill_one_copy = (128+64)*1024  /* bytes */
  };

  /* Do not call hsa_amd_memory_fill when any of the following conditions
     is true. Note that it is always preferred if available and
     before == tail == 0.  */
  if (__builtin_expect (!hsa_fns.hsa_amd_memory_fill_fn, 0)
      || (before && tail && count < always_use_fill)
      || ((before || tail) && count < use_fill_one_copy))
    before = count;

  /* Copy call for alignment - or all data, if condition above is true.  */
  if (before)
    {
      void *data;
      if (before > alloca_size)
	data = malloc (before * sizeof (uint8_t));
      else
	data = alloca (before * sizeof (uint8_t));
      memset (data, val, before);
      status = hsa_fns.hsa_memory_copy_fn (ptr, data, before);
      if (before > alloca_size)
	free (data);
      if (data == 0 || status != HSA_STATUS_SUCCESS)
	goto fail;
      count -= before;
    }

  if (count == 0)
    return true;

  ptr += before;

  uint32_t values = v8 | (v8 << 8) | (v8 << 16) | (v8 << 24);
  status = hsa_fns.hsa_amd_memory_fill_fn (ptr, values, count / 4);
  if (tail && status == HSA_STATUS_SUCCESS)
    {
      ptr += count - tail;
      status = hsa_fns.hsa_memory_copy_fn (ptr, &values, tail);
    }
  if (status == HSA_STATUS_SUCCESS)
    return true;

fail:
  GOMP_PLUGIN_error ("memory set failed");
  return false;
}

void
GOMP_OFFLOAD_interop (struct interop_obj_t *obj, int ord,
		      enum gomp_interop_flag action, bool targetsync,
		      const char *prefer_type)
{
  if ((action == gomp_interop_flag_destroy || action == gomp_interop_flag_use)
      && !obj->stream)
    return;
  if ((action == gomp_interop_flag_destroy || action == gomp_interop_flag_use)
      && obj->fr == omp_ifr_hsa)
    {
      /* Wait until the queue is is empty.   */
      bool is_empty;
      uint64_t read_index, write_index;
      hsa_queue_t *queue = (hsa_queue_t *) obj->stream;
      do
	{
	  read_index = hsa_fns.hsa_queue_load_read_index_relaxed_fn (queue);
	  write_index = hsa_fns.hsa_queue_load_write_index_relaxed_fn (queue);
	  is_empty = (read_index == write_index);
	}
      while (!is_empty);

      if (action == gomp_interop_flag_destroy)
	{
	  hsa_status_t status = hsa_fns.hsa_queue_destroy_fn (queue);
	  if (status != HSA_STATUS_SUCCESS)
	    hsa_fatal ("Error destroying interop hsa_queue_t", status);
	}
      return;
    }
  if (action == gomp_interop_flag_destroy)
    {
      hipError_t err = hip_fns.hipStreamDestroy_fn ((hipStream_t) obj->stream);
      if (err != 0)
	GOMP_PLUGIN_fatal ("Error destroying interop hipStream_t: %d", err);
      return;
    }
  if (action == gomp_interop_flag_use)
    {
      hipError_t err
	= hip_fns.hipStreamSynchronize_fn ((hipStream_t) obj->stream);
      if (err != 0)
	GOMP_PLUGIN_fatal ("Error synchronizing interop hipStream_t: %d", err);
      return;
    }

  bool fr_set = false;

  /* Check for the preferred type; cf. parser in C/C++/Fortran or
     dump_omp_init_prefer_type for the format.
     Accept the first '{...}' block that specifies a 'fr' that we support.
     Currently, no 'attr(...)' are supported.  */
  if (prefer_type)
    while (prefer_type[0] == (char) GOMP_INTEROP_IFR_SEPARATOR)
      {
	/* '{' item block starts.  */
	prefer_type++;
	/* 'fr(...)' block  */
	while (prefer_type[0] != (char) GOMP_INTEROP_IFR_SEPARATOR)
	  {
	    omp_interop_fr_t fr = (omp_interop_fr_t) prefer_type[0];
	    if (fr == omp_ifr_hip)
	      {
		obj->fr = omp_ifr_hip;
		fr_set = true;
	      }
	    if (fr == omp_ifr_hsa)
	      {
		obj->fr = omp_ifr_hsa;
		fr_set = true;
	      }
	    prefer_type++;
	  }
	prefer_type++;
	/* 'attr(...)' block  */
	while (prefer_type[0] != '\0')
	  {
	    /* const char *attr = prefer_type;  */
	    prefer_type += strlen (prefer_type) + 1;
	  }
	prefer_type++;
	/* end of '}'.  */
	if (fr_set)
	  break;
      }

  /* Prefer HIP, use HSA as fallback.  The warning is only printed if GCN_DEBUG
     is set and does not distinguishes between on prefer_type or hip prefer_type
     nor whether a later/lower preference also specifies 'hsa'.
     The assumption is that the user code handles HSA gracefully, but likely
     just by falling back to the host version.  On the other hand, have_hip is
     likely true if HSA is available.  */
  if (!fr_set || obj->fr == omp_ifr_hip)
    {
      bool have_hip = init_hip_runtime_functions ();
      if (have_hip)
	obj->fr = omp_ifr_hip;
      else
	{
	  GCN_WARNING ("interop object requested, using HSA instead of HIP "
		       "as %s could not be loaded", hip_runtime_lib);
	  obj->fr = omp_ifr_hsa;
	}
    }

  _Static_assert (sizeof (uint64_t) == sizeof (hsa_agent_t),
		  "sizeof (uint64_t) == sizeof (hsa_agent_t)");
  struct agent_info *agent = get_agent_info (ord);
  obj->device_data = agent;

  if (targetsync && obj->fr == omp_ifr_hsa)
    {
      hsa_status_t status;
      /* Queue size must be (for GPUs) a power of 2 >= 40, i.e. at least 64 and
	 maximally HSA_AGENT_INFO_QUEUE_MAX_SIZE. Arbitrary choice:  */
      uint32_t queue_size = ASYNC_QUEUE_SIZE;
      status = hsa_fns.hsa_queue_create_fn (agent->id, queue_size,
					    HSA_QUEUE_TYPE_MULTI,
					    NULL, NULL, UINT32_MAX, UINT32_MAX,
					    (hsa_queue_t **) &obj->stream);
      if (status != HSA_STATUS_SUCCESS)
	hsa_fatal ("Error creating interop hsa_queue_t", status);
    }
  else if (targetsync)
    {
      hipError_t err;
      int dev_curr;
      err = hip_fns.hipGetDevice_fn (&dev_curr);
      if (!err && ord != dev_curr)
	err = hip_fns.hipSetDevice_fn (ord);
      if (!err)
	err = hip_fns.hipStreamCreate_fn ((hipStream_t *) &obj->stream);
      if (!err && ord != dev_curr)
	err = hip_fns.hipSetDevice_fn (dev_curr);
      if (err != 0)
	GOMP_PLUGIN_fatal ("Error creating interop hipStream_t: %d", err);
    }
}

intptr_t
GOMP_OFFLOAD_get_interop_int (struct interop_obj_t *obj,
			      omp_interop_property_t property_id,
			      omp_interop_rc_t *ret_code)
{
  if (obj->fr != omp_ifr_hip && obj->fr != omp_ifr_hsa)
    {
      if (ret_code)
	*ret_code = omp_irc_no_value;  /* Hmm. */
      return 0;
    }
  switch (property_id)
    {
    case omp_ipr_fr_id:
      if (ret_code)
	*ret_code = omp_irc_success;
      return obj->fr;
    case omp_ipr_fr_name:
      if (ret_code)
	*ret_code = omp_irc_type_str;
      return 0;
    case omp_ipr_vendor:
      if (ret_code)
	*ret_code = omp_irc_success;
      return 5; /* gnu */
    case omp_ipr_vendor_name:
      if (ret_code)
	*ret_code = omp_irc_type_str;
      return 0;
    case omp_ipr_device_num:
      if (ret_code)
	*ret_code = omp_irc_success;
      return obj->device_num;
    case omp_ipr_platform:
      if (ret_code)
	*ret_code = omp_irc_no_value;
      return 0;
    case omp_ipr_device:
      if (obj->fr == omp_ifr_hsa)
	{
	  if (ret_code)
	    *ret_code = omp_irc_type_ptr;
	  return 0;
	}
	if (ret_code)
	  *ret_code = omp_irc_success;
	return ((struct agent_info *) obj->device_data)->device_id;
    case omp_ipr_device_context:
      if (ret_code && obj->fr == omp_ifr_hsa)
	*ret_code = omp_irc_no_value;
      else if (ret_code)
	*ret_code = omp_irc_type_ptr;
      return 0;
    case omp_ipr_targetsync:
      if (ret_code && !obj->stream)
	*ret_code = omp_irc_no_value;
      else if (ret_code)
	*ret_code = omp_irc_type_ptr;
      return 0;
    default:
      break;
    }
  __builtin_unreachable ();
  return 0;
}

void *
GOMP_OFFLOAD_get_interop_ptr (struct interop_obj_t *obj,
			      omp_interop_property_t property_id,
			      omp_interop_rc_t *ret_code)
{
  if (obj->fr != omp_ifr_hip && obj->fr != omp_ifr_hsa)
    {
      if (ret_code)
	*ret_code = omp_irc_no_value;  /* Hmm. */
      return 0;
    }
  switch (property_id)
    {
    case omp_ipr_fr_id:
      if (ret_code)
	*ret_code = omp_irc_type_int;
      return NULL;
    case omp_ipr_fr_name:
      if (ret_code)
	*ret_code = omp_irc_type_str;
      return NULL;
    case omp_ipr_vendor:
      if (ret_code)
	*ret_code = omp_irc_type_str;
      return NULL;
    case omp_ipr_vendor_name:
      if (ret_code)
	*ret_code = omp_irc_type_str;
      return NULL;
    case omp_ipr_device_num:
      if (ret_code)
	*ret_code = omp_irc_type_int;
      return NULL;
    case omp_ipr_platform:
      if (ret_code)
	*ret_code = omp_irc_no_value;
      return NULL;
    case omp_ipr_device:
      if (obj->fr == omp_ifr_hsa)
	{
	  if (ret_code)
	    *ret_code = omp_irc_success;
	  /* hsa_agent_t is an struct containing a single uint64_t. */
	  return &((struct agent_info *) obj->device_data)->id;
	}
      else
	{
	  if (ret_code)
	    *ret_code = omp_irc_type_int;
	  return NULL;
	}
    case omp_ipr_device_context:
      if (obj->fr == omp_ifr_hsa)
	{
	  if (ret_code)
	    *ret_code = omp_irc_no_value;
	  return NULL;
	}
      else
	{
	  hipCtx_t ctx;
	  int dev_curr;
	  int dev = ((struct agent_info *) obj->device_data)->device_id;
	  hipError_t err;
	  err = hip_fns.hipGetDevice_fn (&dev_curr);
	  if (!err && dev != dev_curr)
	    err = hip_fns.hipSetDevice_fn (dev);
	  if (!err)
	    err = hip_fns.hipCtxGetCurrent_fn (&ctx);
	  if (!err && dev != dev_curr)
	    err = hip_fns.hipSetDevice_fn (dev_curr);
	  if (err)
	    GOMP_PLUGIN_fatal ("Error obtaining hipCtx_t for device %d: %d",
			       obj->device_num, err);
	  if (ret_code)
	    *ret_code = omp_irc_success;
	  return ctx;
	}
    case omp_ipr_targetsync:
      if (!obj->stream)
	{
	  if (ret_code)
	    *ret_code = omp_irc_no_value;
	  return NULL;
	}
      if (ret_code)
	*ret_code = omp_irc_success;
      return obj->stream;
    default:
      break;
    }
  __builtin_unreachable ();
  return NULL;
}

const char *
GOMP_OFFLOAD_get_interop_str (struct interop_obj_t *obj,
			      omp_interop_property_t property_id,
			      omp_interop_rc_t *ret_code)
{
  if (obj->fr != omp_ifr_hip && obj->fr != omp_ifr_hsa)
    {
      if (ret_code)
	*ret_code = omp_irc_no_value;  /* Hmm. */
      return 0;
    }
  switch (property_id)
    {
    case omp_ipr_fr_id:
      if (ret_code)
	*ret_code = omp_irc_type_int;
      return NULL;
    case omp_ipr_fr_name:
      if (ret_code)
	*ret_code = omp_irc_success;
      if (obj->fr == omp_ifr_hip)
	return "hip";
      if (obj->fr == omp_ifr_hsa)
	return "hsa";
    case omp_ipr_vendor:
      if (ret_code)
	*ret_code = omp_irc_type_int;
      return NULL;
    case omp_ipr_vendor_name:
      if (ret_code)
	*ret_code = omp_irc_success;
      return "gnu";
    case omp_ipr_device_num:
      if (ret_code)
	*ret_code = omp_irc_type_int;
      return NULL;
    case omp_ipr_platform:
      if (ret_code)
	*ret_code = omp_irc_no_value;
      return NULL;
    case omp_ipr_device:
      if (ret_code && obj->fr == omp_ifr_hsa)
	*ret_code = omp_irc_type_ptr;
      else if (ret_code)
	*ret_code = omp_irc_type_int;
      return NULL;
    case omp_ipr_device_context:
      if (ret_code && obj->fr == omp_ifr_hsa)
	*ret_code = omp_irc_no_value;
      else if (ret_code)
	*ret_code = omp_irc_type_ptr;
      return NULL;
    case omp_ipr_targetsync:
      if (ret_code && !obj->stream)
	*ret_code = omp_irc_no_value;
      else if (ret_code)
	*ret_code = omp_irc_type_ptr;
      return NULL;
    default:
      break;
    }
  __builtin_unreachable ();
  return 0;
}

const char *
GOMP_OFFLOAD_get_interop_type_desc (struct interop_obj_t *obj,
				    omp_interop_property_t property_id)
{
  _Static_assert (omp_ipr_targetsync == omp_ipr_first,
		  "omp_ipr_targetsync == omp_ipr_first");
  _Static_assert (omp_ipr_platform - omp_ipr_first + 1 == 4,
		  "omp_ipr_platform - omp_ipr_first + 1 == 4");
  static const char *desc_hip[] = {"N/A",		/* platform */
				   "hipDevice_t",	/* device */
				   "hipCtx_t",		/* device_context */
				   "hipStream_t"};	/* targetsync */
  static const char *desc_hsa[] = {"N/A",		/* platform */
				   "hsa_agent_t *",	/* device */
				   "N/A",		/* device_context */
				   "hsa_queue_t *"};	/* targetsync */
  if (obj->fr == omp_ifr_hip)
    return desc_hip[omp_ipr_platform - property_id];
  else
    return desc_hsa[omp_ipr_platform - property_id];
  return NULL;
}

/* }}}  */
/* {{{ 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 (struct gomp_offload_session *session, void *fn_ptr, void **args)
{
  struct agent_info *agent = session->agent;
  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 (session, kernel, 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 (struct gomp_offload_session *session,
			void *tgt_fn,
			void **args,
			void *async_data)
{
  GCN_DEBUG ("GOMP_OFFLOAD_async_run invoked\n");
  struct agent_info *agent = session->agent;
  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);
  if (!agent->omp_async_queue)
    GOMP_PLUGIN_fatal ("Asynchronous queue initialization failed");
  queue_push_launch (agent->omp_async_queue, session, kernel, kla);
  queue_push_callback (agent->omp_async_queue,
		       GOMP_PLUGIN_target_task_completion, async_data);
}

/* Allocate memory suitable for Managed Memory.  */

void *
GOMP_OFFLOAD_managed_alloc (int device, size_t size)
{
  struct agent_info *agent = get_agent_info (device);
  while (1)
    {
      void *result = gomp_simple_alloc (managed_ctx, size);
      if (result)
	return result;

      /* Allocation failed.  Try again if we can create a new heap block.
	 Note: it's possible another thread could get to the new memory
	 first, so the while loop is necessary. */
      if (!managed_heap_create (agent, size))
	return NULL;
    }
}

/* Free memory allocated via GOMP_OFFLOAD_managed_alloc.  */

bool
GOMP_OFFLOAD_managed_free (int device, void *ptr)
{
  gomp_simple_free (managed_ctx, ptr);
  return true;
}

enum accessible {
  UNKNOWN,
  INACCESSIBLE,
  ACCESSIBLE
};

/* Is a host memory address accessible on the given device?
   Returns UNKNOWN if the memory isn't registered, or if it isn't a valid host
   pointer.  */

static enum accessible
host_memory_is_accessible (hsa_agent_t agent, const void *ptr, size_t size)
{
  if (!hsa_fns.hsa_amd_svm_attributes_get_fn)
    return UNKNOWN;

  /* The HSA API doesn't seem to report for the whole range given, so we call
     once for each page the range straddles.  */
  const void *p = ptr;
  size_t remaining = size;
  do
    {
      /* Note: the access query returns in the attribute field.  */
      struct hsa_amd_svm_attribute_pair_s attr = {
	HSA_AMD_SVM_ATTRIB_ACCESS_QUERY, agent.handle
      };
      hsa_status_t status = hsa_fns.hsa_amd_svm_attributes_get_fn ((void*)p,
								   remaining,
								   &attr, 1);
      if (status != HSA_STATUS_SUCCESS)
	/* This happens when the memory isn't registered with ROCr at all.  */
	return UNKNOWN;

      switch (attr.attribute)
	{
	case HSA_AMD_SVM_ATTRIB_AGENT_ACCESSIBLE:
	case HSA_AMD_SVM_ATTRIB_AGENT_ACCESSIBLE_IN_PLACE:
	  break;
	case HSA_AMD_SVM_ATTRIB_AGENT_NO_ACCESS:
	default:
	  return INACCESSIBLE;
	}

      p = (void*)(((uintptr_t)p + 4096) & ~0xfffUL);
      remaining = size - ((uintptr_t)p - (uintptr_t)ptr);
    } while (p < ptr + size);

  /* All pages were accessible.  */
  return ACCESSIBLE;
}

/* Is a device memory address accessible on the given device?
   Returns UNKNOWN if it isn't a valid device address.  Returns INACCESSIBLE if
   the pointer is valid, but not the whole range, or if it refers to the wrong
   device.  */

static enum accessible
device_memory_is_accessible (hsa_agent_t agent, const void *ptr, size_t size)
{
  if (!hsa_fns.hsa_amd_pointer_info_fn)
    return UNKNOWN;

  hsa_amd_pointer_info_t info;
  uint32_t nagents;
  hsa_agent_t *agents;
  info.size = sizeof (hsa_amd_pointer_info_t);

  hsa_status_t status = hsa_fns.hsa_amd_pointer_info_fn (ptr, &info, NULL,
							 &nagents, &agents); 
  if (status != HSA_STATUS_SUCCESS
      || info.type == HSA_EXT_POINTER_TYPE_UNKNOWN)
    return UNKNOWN;

  if (agent.handle == info.agentOwner.handle)
    return (info.sizeInBytes >= size ? ACCESSIBLE : INACCESSIBLE);

  for (unsigned i = 0; i < nagents; i++)
    {
      if (agent.handle == agents[0].handle)
	return (info.sizeInBytes >= size ? ACCESSIBLE : INACCESSIBLE); 
    }

  return INACCESSIBLE;
}

/* Backend implementation for omp_target_is_accessible.  */

int
GOMP_OFFLOAD_is_accessible_ptr (int device, const void *ptr, size_t size)
{
  if (!init_hsa_context (false)
      || device < 0 || device > hsa_context.agent_count)
    return 0;

  struct agent_info *agent = get_agent_info (device);

  enum accessible result;
  result = host_memory_is_accessible (agent->id, ptr, size);
  if (result == UNKNOWN)
    result = device_memory_is_accessible (agent->id, ptr, size);
  return result == ACCESSIBLE;
}

/* }}} */
/* {{{ OpenACC Plugin API  */

/* Run a synchronous OpenACC kernel.  The device number is inferred from the
   already-loaded KERNEL.  */

void
GOMP_OFFLOAD_openacc_exec (struct gomp_offload_session *session,
			   void (*fn_ptr) (void *),
			   size_t mapnum __attribute__((unused)),
			   void **hostaddrs __attribute__((unused)),
			   unsigned *dims, void *targ_mem_desc)
{
  struct kernel_info *kernel = (struct kernel_info *) fn_ptr;

  gcn_exec (kernel, session, dims, targ_mem_desc, false, NULL);
}

/* Run an asynchronous OpenACC kernel on the specified queue.  */

void
GOMP_OFFLOAD_openacc_async_exec (struct gomp_offload_session *session,
				 void (*fn_ptr) (void *),
				 size_t mapnum __attribute__((unused)),
				 void **hostaddrs __attribute__((unused)),
				 unsigned *dims, void *targ_mem_desc,
				 struct goacc_asyncqueue *aq)
{
  struct kernel_info *kernel = (struct kernel_info *) fn_ptr;

  gcn_exec (kernel, session, 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 NULL;
    }
  if (pthread_cond_init (&aq->queue_cond_in, NULL))
    {
      GOMP_PLUGIN_error ("Failed to initialize a GCN agent queue cond");
      return NULL;
    }
  if (pthread_cond_init (&aq->queue_cond_out, NULL))
    {
      GOMP_PLUGIN_error ("Failed to initialize a GCN agent queue cond");
      return NULL;
    }

  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 existing 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.
   (Also handles dev2host and dev2dev.)  */

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);
  queue_push_copy (aq, dst, src, n);
  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)
{
  return GOMP_OFFLOAD_openacc_async_host2dev (device, dst, src, n, aq);
}

/* Queue up an asynchronous data copy from DEVICE to DEVICE.  */

bool
GOMP_OFFLOAD_openacc_async_dev2dev (int device, void *dst, const void *src,
				    size_t n, struct goacc_asyncqueue *aq)
{
  return GOMP_OFFLOAD_openacc_async_host2dev (device, dst, src, n, aq);
}

union goacc_property_value
GOMP_OFFLOAD_openacc_get_property (int device, enum goacc_property prop)
{
  struct agent_info *agent = get_agent_info (device);

  union goacc_property_value propval = { .val = 0 };

  switch (prop)
    {
    case GOACC_PROPERTY_FREE_MEMORY:
      /* Not supported. */
      break;
    case GOACC_PROPERTY_MEMORY:
      {
	size_t size;
	hsa_region_t region = agent->data_region;
	hsa_status_t status =
	  hsa_fns.hsa_region_get_info_fn (region, HSA_REGION_INFO_SIZE, &size);
	if (status == HSA_STATUS_SUCCESS)
	  propval.val = size;
	break;
      }
    case GOACC_PROPERTY_NAME:
      propval.ptr = agent->name;
      break;
    case GOACC_PROPERTY_VENDOR:
      propval.ptr = agent->vendor_name;
      break;
    case GOACC_PROPERTY_DRIVER:
      propval.ptr = hsa_context.driver_version_s;
      break;
    }

  return propval;
}

/* 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);
}

#ifdef USE_HASHTAB_LOOKUP_FOR_INDIRECT
  #include "build-target-indirect-htab.h"
#endif

/* }}} */