Removal of HSA offloading from gcc and libgomp

This patch removes the generation of HSAIL from the compiler, the HSA
offloading plugin from libgomp and the associated testsuite tests and
infrastructure bits from the respective testsuites.

Apart from removal of the obvious files, I removed bits that I found
by searching for HSA related terms and by re-tracing my steps and
looking at the patches that introduced HSA in the first place.  I did
not remove everything these patches brought in, for example:

  - the mechanism to pass offload-target specific info from the application to
    the offloading plugin - but the same mechanism is also used to
    communicate number of teams and the thread limit to all offload targets.

  - run_func hook in gomp_device_descr stays too, although now it is
    not used.  If some future offload target would like the ability to
    refuse to offload some functions, it can use it.  It is easy to
    remove as a follow-up if it is considered clutter, though.

  - configure options --with-hsa-runtime=PATH, -with-hsa-runtime-include=PATH
    and --with-hsa-runtime-lib=PATH rmeain because GCN uses them too.

  - Surprisingly, GOMP_TARGET_ARG_HSA_KERNEL_ATTRIBUTES (a constant
    from gomp-constants.h) appears in the source of the amdgcn libgomp
    plugin, although I tend to think that code path is not ever used
    and this patch certainly removes it from the compiler.
    Nevertheless, it seems it has potential value beyond HSAIL and so
    I've kept it, it can of course always be easily removed in the
    future of GCN folk abandon it too.

  - I assume constants OFFLOAD_TARGET_TYPE_HSA and GOMP_DEVICE_HSA
    need to stay indefinitely too just so that no future offload
    target picks that number.

  - I have kept dg-require-effective-target
    offload_device_nonshared_as requirement of thests which have it.

It is quite probable I missed some small HSA artifacts but those
should be easy to remove later as we find them.

include/ChangeLog:

2020-07-24  Martin Jambor  <mjambor@suse.cz>

	* gomp-constants.h (GOMP_VERSION_HSA): Remove.

gcc/ChangeLog:

2020-07-24  Martin Jambor  <mjambor@suse.cz>

	* hsa-brig-format.h: Moved to brig/brigfrontend.
	* hsa-brig.c: Removed.
	* hsa-builtins.def: Likewise.
	* hsa-common.c: Likewise.
	* hsa-common.h: Likewise.
	* hsa-dump.c: Likewise.
	* hsa-gen.c: Likewise.
	* hsa-regalloc.c: Likewise.
	* ipa-hsa.c: Likewise.
	* omp-grid.c: Likewise.
	* omp-grid.h: Likewise.
	* Makefile.in (BUILTINS_DEF): Remove hsa-builtins.def.
	(OBJS): Remove hsa-common.o, hsa-gen.o, hsa-regalloc.o, hsa-brig.o,
	hsa-dump.o, ipa-hsa.c and omp-grid.o.
	(GTFILES): Removed hsa-common.c and omp-expand.c.
	* builtins.def: Remove processing of hsa-builtins.def.
	(DEF_HSA_BUILTIN): Remove.
	* common.opt (flag_disable_hsa): Remove.
	(-Whsa): Ignore.
	* config.in (ENABLE_HSA): Removed.
	* configure.ac: Removed handling configuration for hsa offloading.
	(ENABLE_HSA): Removed.
	* configure: Regenerated.
	* doc/install.texi (--enable-offload-targets): Remove hsa from the
	example.
	(--with-hsa-runtime): Reword to reference any HSA run-time, not
	specifically HSA offloading.
	* doc/invoke.texi (Option Summary): Remove -Whsa.
	(Warning Options): Likewise.
	(Optimize Options): Remove hsa-gen-debug-stores.
	* doc/passes.texi (Regular IPA passes): Remove section on IPA HSA
	pass.
	* gimple-low.c (lower_stmt): Remove GIMPLE_OMP_GRID_BODY case.
	* gimple-pretty-print.c (dump_gimple_omp_for): Likewise.
	(dump_gimple_omp_block): Likewise.
	(pp_gimple_stmt_1): Likewise.
	* gimple-walk.c (walk_gimple_stmt): Likewise.
	* gimple.c (gimple_build_omp_grid_body): Removed function.
	(gimple_copy): Remove GIMPLE_OMP_GRID_BODY case.
	* gimple.def (GIMPLE_OMP_GRID_BODY): Removed.
	* gimple.h (gf_mask): Removed GF_OMP_PARALLEL_GRID_PHONY,
	OMP_FOR_KIND_GRID_LOOP, GF_OMP_FOR_GRID_PHONY,
	GF_OMP_FOR_GRID_INTRA_GROUP, GF_OMP_FOR_GRID_GROUP_ITER and
	GF_OMP_TEAMS_GRID_PHONY.  Renumbered GF_OMP_FOR_KIND_SIMD and
	GF_OMP_TEAMS_HOST.
	(gimple_build_omp_grid_body): Removed declaration.
	(gimple_has_substatements): Remove GIMPLE_OMP_GRID_BODY case.
	(gimple_omp_for_grid_phony): Removed.
	(gimple_omp_for_set_grid_phony): Likewise.
	(gimple_omp_for_grid_intra_group): Likewise.
	(gimple_omp_for_grid_intra_group): Likewise.
	(gimple_omp_for_grid_group_iter): Likewise.
	(gimple_omp_for_set_grid_group_iter): Likewise.
	(gimple_omp_parallel_grid_phony): Likewise.
	(gimple_omp_parallel_set_grid_phony): Likewise.
	(gimple_omp_teams_grid_phony): Likewise.
	(gimple_omp_teams_set_grid_phony): Likewise.
	(CASE_GIMPLE_OMP): Remove GIMPLE_OMP_GRID_BODY case.
	* lto-section-in.c (lto_section_name): Removed hsa.
	* lto-streamer.h (lto_section_type): Removed LTO_section_ipa_hsa.
	* lto-wrapper.c (compile_images_for_offload_targets): Remove special
	handling of hsa.
	* omp-expand.c: Do not include hsa-common.h and gt-omp-expand.h.
	(parallel_needs_hsa_kernel_p): Removed.
	(grid_launch_attributes_trees): Likewise.
	(grid_launch_attributes_trees): Likewise.
	(grid_create_kernel_launch_attr_types): Likewise.
	(grid_insert_store_range_dim): Likewise.
	(grid_get_kernel_launch_attributes): Likewise.
	(get_target_arguments): Remove code passing HSA grid sizes.
	(grid_expand_omp_for_loop): Remove.
	(grid_arg_decl_map): Likewise.
	(grid_remap_kernel_arg_accesses): Likewise.
	(grid_expand_target_grid_body): Likewise.
	(expand_omp): Remove call to grid_expand_target_grid_body.
	(omp_make_gimple_edges): Remove GIMPLE_OMP_GRID_BODY case.
	* omp-general.c: Do not include hsa-common.h.
	(omp_maybe_offloaded): Do not check for HSA offloading.
	(omp_context_selector_matches): Likewise.
	* omp-low.c: Do not include hsa-common.h and omp-grid.h.
	(build_outer_var_ref): Remove handling of GIMPLE_OMP_GRID_BODY.
	(scan_sharing_clauses): Remove handling of OMP_CLAUSE__GRIDDIM_.
	(scan_omp_parallel): Remove handling of the phoney variant.
	(check_omp_nesting_restrictions): Remove handling of
	GIMPLE_OMP_GRID_BODY and GF_OMP_FOR_KIND_GRID_LOOP.
	(scan_omp_1_stmt): Remove handling of GIMPLE_OMP_GRID_BODY.
	(lower_omp_for_lastprivate): Remove handling of gridified loops.
	(lower_omp_for): Remove phony loop handling.
	(lower_omp_taskreg): Remove phony construct handling.
	(lower_omp_teams): Likewise.
	(lower_omp_grid_body): Removed.
	(lower_omp_1): Remove GIMPLE_OMP_GRID_BODY case.
	(execute_lower_omp): Do not call omp_grid_gridify_all_targets.
	* opts.c (common_handle_option): Do not handle hsa when processing
	OPT_foffload_.
	* params.opt (hsa-gen-debug-stores): Remove.
	* passes.def: Remove pass_ipa_hsa and pass_gen_hsail.
	* timevar.def: Remove TV_IPA_HSA.
	* toplev.c: Do not include hsa-common.h.
	(compile_file): Do not call hsa_output_brig.
	* tree-core.h (enum omp_clause_code): Remove OMP_CLAUSE__GRIDDIM_.
	(tree_omp_clause): Remove union field dimension.
	* tree-nested.c (convert_nonlocal_omp_clauses): Remove the
	OMP_CLAUSE__GRIDDIM_ case.
	(convert_local_omp_clauses): Likewise.
	* tree-pass.h (make_pass_gen_hsail): Remove declaration.
	(make_pass_ipa_hsa): Likewise.
	* tree-pretty-print.c (dump_omp_clause): Remove GIMPLE_OMP_GRID_BODY
	case.
	* tree.c (omp_clause_num_ops): Remove the element corresponding to
	OMP_CLAUSE__GRIDDIM_.
	(omp_clause_code_name): Likewise.
	(walk_tree_1): Remove GIMPLE_OMP_GRID_BODY case.
	* tree.h (OMP_CLAUSE__GRIDDIM__DIMENSION): Remove.
	(OMP_CLAUSE__GRIDDIM__SIZE): Likewise.
	(OMP_CLAUSE__GRIDDIM__GROUP): Likewise.

gcc/fortran/ChangeLog:

2020-07-24  Martin Jambor  <mjambor@suse.cz>

	* f95-lang.c (gfc_init_builtin_functions): Remove processing of
	hsa-builtins.def.

gcc/brig/ChangeLog:

2020-07-24  Martin Jambor  <mjambor@suse.cz>

	* brigfrontend/brig-util.h (hsa_type_packed_p): Declared.
	* brigfrontend/brig-util.cc (hsa_type_packed_p): Moved here from
	removed gcc/hsa-common.c.

libgomp/ChangeLog:

2020-07-24  Martin Jambor  <mjambor@suse.cz>

	* plugin/Makefrag.am: Remove configuration of HSA plugin.
	* aclocal.m4: Regenerated.
	* Makefile.in: Regenerated.
	* config.h.in: Regenerated.
	* configure: Regenerated.
	* plugin/configfrag.ac: Likewise.
	* plugin/hsa_ext_finalize.h: Removed.
	* plugin/plugin-hsa.c: Likewise.
	* testsuite/Makefile.in: Regenerated.
	* testsuite/lib/libgomp.exp
	(offload_target_to_openacc_device_type): Remove hsa case.
	(check_effective_target_hsa_offloading_selected_nocache): Removed
	(check_effective_target_hsa_offloading_selected): Likewise.
	(libgomp_init): Do not add -Wno-hsa to additional_flags.
	* testsuite/libgomp.hsa.c/alloca-1.c: Removed test.
	* testsuite/libgomp.hsa.c/bitfield-1.c: Likewise.
	* testsuite/libgomp.hsa.c/bits-insns.c: Likewise.
	* testsuite/libgomp.hsa.c/builtins-1.c: Likewise.
	* testsuite/libgomp.hsa.c/c.exp: Likewise.
	* testsuite/libgomp.hsa.c/complex-1.c: Likewise.
	* testsuite/libgomp.hsa.c/complex-align-2.c: Likewise.
	* testsuite/libgomp.hsa.c/formal-actual-args-1.c: Likewise.
	* testsuite/libgomp.hsa.c/function-call-1.c: Likewise.
	* testsuite/libgomp.hsa.c/get-level-1.c: Likewise.
	* testsuite/libgomp.hsa.c/gridify-1.c: Likewise.
	* testsuite/libgomp.hsa.c/gridify-2.c: Likewise.
	* testsuite/libgomp.hsa.c/gridify-3.c: Likewise.
	* testsuite/libgomp.hsa.c/gridify-4.c: Likewise.
	* testsuite/libgomp.hsa.c/memory-operations-1.c: Likewise.
	* testsuite/libgomp.hsa.c/pr69568.c: Likewise.
	* testsuite/libgomp.hsa.c/pr82416.c: Likewise.
	* testsuite/libgomp.hsa.c/rotate-1.c: Likewise.
	* testsuite/libgomp.hsa.c/staticvar.c: Likewise.
	* testsuite/libgomp.hsa.c/switch-1.c: Likewise.
	* testsuite/libgomp.hsa.c/switch-branch-1.c: Likewise.
	* testsuite/libgomp.hsa.c/switch-sbr-2.c: Likewise.
	* testsuite/libgomp.hsa.c/tiling-1.c: Likewise.
	* testsuite/libgomp.hsa.c/tiling-2.c: Likewise.

gcc/testsuite/ChangeLog:

2020-07-24  Martin Jambor  <mjambor@suse.cz>

	* lib/target-supports.exp (check_effective_target_offload_hsa):
	Removed.
	* c-c++-common/gomp/gridify-1.c: Removed test.
	* c-c++-common/gomp/gridify-2.c: Likewise.
	* c-c++-common/gomp/gridify-3.c: Likewise.
	* c-c++-common/gomp/hsa-indirect-call-1.c: Likewise.
	* gfortran.dg/gomp/gridify-1.f90: Likewise.
	* gcc.dg/gomp/gomp.exp: Do not pass -Wno-hsa to tests.
	* g++.dg/gomp/gomp.exp: Likewise.
	* gfortran.dg/gomp/gomp.exp: Likewise.

4 files changed, 1 insertions, 2210 deletions

diff --git a/libgomp/plugin/Makefrag.am b/libgomp/plugin/Makefrag.am
index 7924959..54301f8 100644
--- a/libgomp/plugin/Makefrag.am
+++ b/libgomp/plugin/Makefrag.am
@@ -39,20 +39,6 @@ libgomp_plugin_nvptx_la_LIBADD = libgomp.la $(PLUGIN_NVPTX_LIBS)
 libgomp_plugin_nvptx_la_LIBTOOLFLAGS = --tag=disable-static
 endif
 
-if PLUGIN_HSA
-# Heterogenous Systems Architecture plugin
-libgomp_plugin_hsa_version_info = -version-info $(libtool_VERSION)
-toolexeclib_LTLIBRARIES += libgomp-plugin-hsa.la
-libgomp_plugin_hsa_la_SOURCES = plugin/plugin-hsa.c
-libgomp_plugin_hsa_la_CPPFLAGS = $(AM_CPPFLAGS) $(PLUGIN_HSA_CPPFLAGS) \
-	-D_GNU_SOURCE
-libgomp_plugin_hsa_la_LDFLAGS = $(libgomp_plugin_hsa_version_info) \
-	$(lt_host_flags)
-libgomp_plugin_hsa_la_LDFLAGS += $(PLUGIN_HSA_LDFLAGS)
-libgomp_plugin_hsa_la_LIBADD = libgomp.la $(PLUGIN_HSA_LIBS)
-libgomp_plugin_hsa_la_LIBTOOLFLAGS = --tag=disable-static
-endif
-
 if PLUGIN_GCN
 # AMD GCN plugin
 libgomp_plugin_gcn_version_info = -version-info $(libtool_VERSION)
diff --git a/libgomp/plugin/configfrag.ac b/libgomp/plugin/configfrag.ac
index fc91702..1403008 100644
--- a/libgomp/plugin/configfrag.ac
+++ b/libgomp/plugin/configfrag.ac
@@ -128,15 +128,6 @@ if test "x$HSA_RUNTIME_LIB" != x; then
   HSA_RUNTIME_LDFLAGS=-L$HSA_RUNTIME_LIB
 fi
 
-PLUGIN_HSA=0
-PLUGIN_HSA_CPPFLAGS=
-PLUGIN_HSA_LDFLAGS=
-PLUGIN_HSA_LIBS=
-AC_SUBST(PLUGIN_HSA)
-AC_SUBST(PLUGIN_HSA_CPPFLAGS)
-AC_SUBST(PLUGIN_HSA_LDFLAGS)
-AC_SUBST(PLUGIN_HSA_LIBS)
-
 PLUGIN_GCN=0
 PLUGIN_GCN_CPPFLAGS=
 PLUGIN_GCN_LDFLAGS=
@@ -207,45 +198,6 @@ if test x"$enable_offload_targets" != x; then
 	  ;;
 	esac
 	;;
-      hsa*)
-	case "${target}" in
-	  x86_64-*-*)
-	    case " ${CC} ${CFLAGS} " in
-	      *" -m32 "*|*" -mx32 "*)
-	        PLUGIN_HSA=0
-		;;
-	      *)
-		tgt_plugin=hsa
-	        PLUGIN_HSA=$tgt
-	        PLUGIN_HSA_CPPFLAGS=$HSA_RUNTIME_CPPFLAGS
-	        PLUGIN_HSA_LDFLAGS="$HSA_RUNTIME_LDFLAGS"
-	        PLUGIN_HSA_LIBS="-ldl"
-
-	        PLUGIN_HSA_save_CPPFLAGS=$CPPFLAGS
-	        CPPFLAGS="$PLUGIN_HSA_CPPFLAGS $CPPFLAGS"
-	        PLUGIN_HSA_save_LDFLAGS=$LDFLAGS
-	        LDFLAGS="$PLUGIN_HSA_LDFLAGS $LDFLAGS"
-	        PLUGIN_HSA_save_LIBS=$LIBS
-	        LIBS="$PLUGIN_HSA_LIBS $LIBS"
-
-	        PLUGIN_HSA=1
-	        CPPFLAGS=$PLUGIN_HSA_save_CPPFLAGS
-	        LDFLAGS=$PLUGIN_HSA_save_LDFLAGS
-	        LIBS=$PLUGIN_HSA_save_LIBS
-	        case $PLUGIN_HSA in
-		  hsa*)
-	            HSA_PLUGIN=0
-	            AC_MSG_ERROR([HSA run-time package required for HSA support])
-	            ;;
-	        esac
-		;;
-	      esac
-    	    ;;
-	  *-*-*)
-	    PLUGIN_HSA=0
-            ;;
-        esac
-        ;;
 
       amdgcn*)
 	case "${target}" in
@@ -285,10 +237,7 @@ if test x"$enable_offload_targets" != x; then
       offload_targets=$offload_targets,$tgt
     fi
     # Configure additional search paths.
-    if test "$tgt_plugin" = hsa; then
-      # Offloading compilation is all handled by the target compiler.
-      :
-    elif test x"$tgt_dir" != x; then
+    if test x"$tgt_dir" != x; then
       offload_additional_options="$offload_additional_options -B$tgt_dir/libexec/gcc/\$(target_alias)/\$(gcc_version) -B$tgt_dir/bin"
       offload_additional_lib_paths="$offload_additional_lib_paths:$tgt_dir/lib64:$tgt_dir/lib:$tgt_dir/lib32"
     else
@@ -304,9 +253,6 @@ AC_DEFINE_UNQUOTED([PLUGIN_NVPTX], [$PLUGIN_NVPTX],
   [Define to 1 if the NVIDIA plugin is built, 0 if not.])
 AC_DEFINE_UNQUOTED([PLUGIN_NVPTX_DYNAMIC], [$PLUGIN_NVPTX_DYNAMIC],
   [Define to 1 if the NVIDIA plugin should dlopen libcuda.so.1, 0 if it should be linked against it.])
-AM_CONDITIONAL([PLUGIN_HSA], [test $PLUGIN_HSA = 1])
-AC_DEFINE_UNQUOTED([PLUGIN_HSA], [$PLUGIN_HSA],
-  [Define to 1 if the HSA plugin is built, 0 if not.])
 AM_CONDITIONAL([PLUGIN_GCN], [test $PLUGIN_GCN = 1])
 AC_DEFINE_UNQUOTED([PLUGIN_GCN], [$PLUGIN_GCN],
   [Define to 1 if the GCN plugin is built, 0 if not.])
diff --git a/libgomp/plugin/hsa_ext_finalize.h b/libgomp/plugin/hsa_ext_finalize.h
deleted file mode 100644
index 39284bd..0000000
--- a/libgomp/plugin/hsa_ext_finalize.h
+++ /dev/null
@@ -1,270 +0,0 @@
-/* HSA Extensions API 1.0.1 representation description.
-   Copyright (C) 2016-2020 Free Software Foundation, Inc.
-
-This file is part of GCC.
-
-GCC is free software; you can redistribute it and/or modify
-it under the terms of the GNU General Public License as published by
-the Free Software Foundation; either version 3, or (at your option)
-any later version.
-
-GCC is distributed in the hope that it will be useful,
-but WITHOUT ANY WARRANTY; without even the implied warranty of
-MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
-GNU General Public License for more details.
-
-Under Section 7 of GPL version 3, you are granted additional
-permissions described in the GCC Runtime Library Exception, version
-3.1, as published by the Free Software Foundation.
-
-You should have received a copy of the GNU General Public License and
-a copy of the GCC Runtime Library Exception along with this program;
-see the files COPYING3 and COPYING.RUNTIME respectively.  If not, see
-<http://www.gnu.org/licenses/>.
-
-The contents of the file was created by extracting data structures, enum,
-typedef and other definitions from HSA Runtime Programmer’s Reference Manual
-Version 1.0 (http://www.hsafoundation.com/standards/).
-
-HTML version is provided on the following link:
-http://www.hsafoundation.com/html/Content/Runtime/Topics/Runtime_title_page.htm
-*/
-
-
-#ifndef _HSA_EXT_FINALIZE_H
-#define _HSA_EXT_FINALIZE_H 1
-
-struct BrigModuleHeader;
-typedef struct BrigModuleHeader *BrigModule_t;
-
-typedef enum {
-  HSA_EXT_IMAGE_GEOMETRY_1D = 0,
-  HSA_EXT_IMAGE_GEOMETRY_2D = 1,
-  HSA_EXT_IMAGE_GEOMETRY_3D = 2,
-  HSA_EXT_IMAGE_GEOMETRY_1DA = 3,
-  HSA_EXT_IMAGE_GEOMETRY_2DA = 4,
-  HSA_EXT_IMAGE_GEOMETRY_1DB = 5,
-  HSA_EXT_IMAGE_GEOMETRY_2DDEPTH = 6,
-  HSA_EXT_IMAGE_GEOMETRY_2DADEPTH = 7
-} hsa_ext_image_geometry_t;
-
-typedef enum {
-  HSA_EXT_IMAGE_CHANNEL_TYPE_SNORM_INT8 = 0,
-  HSA_EXT_IMAGE_CHANNEL_TYPE_SNORM_INT16 = 1,
-  HSA_EXT_IMAGE_CHANNEL_TYPE_UNORM_INT8 = 2,
-  HSA_EXT_IMAGE_CHANNEL_TYPE_UNORM_INT16 = 3,
-  HSA_EXT_IMAGE_CHANNEL_TYPE_UNORM_INT24 = 4,
-  HSA_EXT_IMAGE_CHANNEL_TYPE_UNORM_SHORT_555 = 5,
-  HSA_EXT_IMAGE_CHANNEL_TYPE_UNORM_SHORT_565 = 6,
-  HSA_EXT_IMAGE_CHANNEL_TYPE_UNORM_SHORT_101010 = 7,
-  HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT8 = 8,
-  HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT16 = 9,
-  HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT32 = 10,
-  HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT8 = 11,
-  HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT16 = 12,
-  HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT32 = 13,
-  HSA_EXT_IMAGE_CHANNEL_TYPE_HALF_FLOAT = 14,
-  HSA_EXT_IMAGE_CHANNEL_TYPE_FLOAT = 15
-} hsa_ext_image_channel_type_t;
-
-typedef enum {
-  HSA_EXT_IMAGE_CHANNEL_ORDER_A = 0,
-  HSA_EXT_IMAGE_CHANNEL_ORDER_R = 1,
-  HSA_EXT_IMAGE_CHANNEL_ORDER_RX = 2,
-  HSA_EXT_IMAGE_CHANNEL_ORDER_RG = 3,
-  HSA_EXT_IMAGE_CHANNEL_ORDER_RGX = 4,
-  HSA_EXT_IMAGE_CHANNEL_ORDER_RA = 5,
-  HSA_EXT_IMAGE_CHANNEL_ORDER_RGB = 6,
-  HSA_EXT_IMAGE_CHANNEL_ORDER_RGBX = 7,
-  HSA_EXT_IMAGE_CHANNEL_ORDER_RGBA = 8,
-  HSA_EXT_IMAGE_CHANNEL_ORDER_BGRA = 9,
-  HSA_EXT_IMAGE_CHANNEL_ORDER_ARGB = 10,
-  HSA_EXT_IMAGE_CHANNEL_ORDER_ABGR = 11,
-  HSA_EXT_IMAGE_CHANNEL_ORDER_SRGB = 12,
-  HSA_EXT_IMAGE_CHANNEL_ORDER_SRGBX = 13,
-  HSA_EXT_IMAGE_CHANNEL_ORDER_SRGBA = 14,
-  HSA_EXT_IMAGE_CHANNEL_ORDER_SBGRA = 15,
-  HSA_EXT_IMAGE_CHANNEL_ORDER_INTENSITY = 16,
-  HSA_EXT_IMAGE_CHANNEL_ORDER_LUMINANCE = 17,
-  HSA_EXT_IMAGE_CHANNEL_ORDER_DEPTH = 18,
-  HSA_EXT_IMAGE_CHANNEL_ORDER_DEPTH_STENCIL = 19
-} hsa_ext_image_channel_order_t;
-
-typedef struct hsa_ext_image_format_s
-{
-  hsa_ext_image_channel_type_t channel_type;
-  hsa_ext_image_channel_order_t channel_order;
-} hsa_ext_image_format_t;
-
-typedef struct hsa_ext_sampler_s
-{
-  uint64_t handle;
-} hsa_ext_sampler_t;
-typedef struct hsa_ext_image_data_info_s
-{
-  size_t size;
-  size_t alignment;
-} hsa_ext_image_data_info_t;
-typedef enum {
-  HSA_EXT_SAMPLER_ADDRESSING_MODE_UNDEFINED = 0,
-  HSA_EXT_SAMPLER_ADDRESSING_MODE_CLAMP_TO_EDGE = 1,
-  HSA_EXT_SAMPLER_ADDRESSING_MODE_CLAMP_TO_BORDER = 2,
-  HSA_EXT_SAMPLER_ADDRESSING_MODE_REPEAT = 3,
-  HSA_EXT_SAMPLER_ADDRESSING_MODE_MIRRORED_REPEAT = 4
-} hsa_ext_sampler_addressing_mode_t;
-typedef struct hsa_ext_image_s
-{
-  uint64_t handle;
-} hsa_ext_image_t;
-typedef enum {
-  HSA_EXT_IMAGE_CAPABILITY_NOT_SUPPORTED = 0x0,
-  HSA_EXT_IMAGE_CAPABILITY_READ_ONLY = 0x1,
-  HSA_EXT_IMAGE_CAPABILITY_WRITE_ONLY = 0x2,
-  HSA_EXT_IMAGE_CAPABILITY_READ_WRITE = 0x4,
-  HSA_EXT_IMAGE_CAPABILITY_READ_MODIFY_WRITE = 0x8,
-  HSA_EXT_IMAGE_CAPABILITY_ACCESS_INVARIANT_DATA_LAYOUT = 0x10
-} hsa_ext_image_capability_t;
-typedef struct hsa_ext_control_directives_s
-{
-  uint64_t control_directives_mask;
-  uint16_t break_exceptions_mask;
-  uint16_t detect_exceptions_mask;
-  uint32_t max_dynamic_group_size;
-  uint64_t max_flat_grid_size;
-  uint32_t max_flat_workgroup_size;
-  uint32_t reserved1;
-  uint64_t required_grid_size[3];
-  hsa_dim3_t required_workgroup_size;
-  uint8_t required_dim;
-  uint8_t reserved2[75];
-} hsa_ext_control_directives_t;
-typedef enum {
-  HSA_EXT_SAMPLER_FILTER_MODE_NEAREST = 0,
-  HSA_EXT_SAMPLER_FILTER_MODE_LINEAR = 1
-} hsa_ext_sampler_filter_mode_t;
-
-typedef enum {
-  HSA_EXT_SAMPLER_COORDINATE_MODE_UNNORMALIZED = 0,
-  HSA_EXT_SAMPLER_COORDINATE_MODE_NORMALIZED = 1
-} hsa_ext_sampler_coordinate_mode_t;
-typedef enum {
-  HSA_EXT_FINALIZER_CALL_CONVENTION_AUTO = -1
-} hsa_ext_finalizer_call_convention_t;
-typedef struct hsa_ext_program_s
-{
-  uint64_t handle;
-} hsa_ext_program_t;
-typedef struct hsa_ext_image_descriptor_s
-{
-  hsa_ext_image_geometry_t geometry;
-  size_t width;
-  size_t height;
-  size_t depth;
-  size_t array_size;
-  hsa_ext_image_format_t format;
-} hsa_ext_image_descriptor_t;
-typedef enum {
-  HSA_EXT_PROGRAM_INFO_MACHINE_MODEL = 0,
-  HSA_EXT_PROGRAM_INFO_PROFILE = 1,
-  HSA_EXT_PROGRAM_INFO_DEFAULT_FLOAT_ROUNDING_MODE = 2
-} hsa_ext_program_info_t;
-typedef BrigModule_t hsa_ext_module_t;
-typedef struct hsa_ext_sampler_descriptor_s
-{
-  hsa_ext_sampler_coordinate_mode_t coordinate_mode;
-  hsa_ext_sampler_filter_mode_t filter_mode;
-  hsa_ext_sampler_addressing_mode_t address_mode;
-} hsa_ext_sampler_descriptor_t;
-
-typedef struct hsa_ext_image_region_s
-{
-  hsa_dim3_t offset;
-  hsa_dim3_t range;
-} hsa_ext_image_region_t;
-hsa_status_t hsa_ext_image_export (hsa_agent_t agent, hsa_ext_image_t src_image,
-				   void *dst_memory, size_t dst_row_pitch,
-				   size_t dst_slice_pitch,
-				   const hsa_ext_image_region_t *image_region);
-hsa_status_t hsa_ext_program_add_module (hsa_ext_program_t program,
-					 hsa_ext_module_t module);
-hsa_status_t hsa_ext_program_iterate_modules (
-  hsa_ext_program_t program,
-  hsa_status_t (*callback) (hsa_ext_program_t program, hsa_ext_module_t module,
-			    void *data),
-  void *data);
-hsa_status_t hsa_ext_program_create (
-  hsa_machine_model_t machine_model, hsa_profile_t profile,
-  hsa_default_float_rounding_mode_t default_float_rounding_mode,
-  const char *options, hsa_ext_program_t *program);
-hsa_status_t
-hsa_ext_image_data_get_info (hsa_agent_t agent,
-			     const hsa_ext_image_descriptor_t *image_descriptor,
-			     hsa_access_permission_t access_permission,
-			     hsa_ext_image_data_info_t *image_data_info);
-
-hsa_status_t hsa_ext_image_import (hsa_agent_t agent, const void *src_memory,
-				   size_t src_row_pitch, size_t src_slice_pitch,
-				   hsa_ext_image_t dst_image,
-				   const hsa_ext_image_region_t *image_region);
-hsa_status_t hsa_ext_program_get_info (hsa_ext_program_t program,
-				       hsa_ext_program_info_t attribute,
-				       void *value);
-enum
-{
-  HSA_EXT_STATUS_ERROR_IMAGE_FORMAT_UNSUPPORTED = 0x3000,
-  HSA_EXT_STATUS_ERROR_IMAGE_SIZE_UNSUPPORTED = 0x3001
-};
-hsa_status_t hsa_ext_image_destroy (hsa_agent_t agent, hsa_ext_image_t image);
-hsa_status_t hsa_ext_image_get_capability (
-  hsa_agent_t agent, hsa_ext_image_geometry_t geometry,
-  const hsa_ext_image_format_t *image_format, uint32_t *capability_mask);
-enum
-{
-  HSA_EXT_STATUS_ERROR_INVALID_PROGRAM = 0x2000,
-  HSA_EXT_STATUS_ERROR_INVALID_MODULE = 0x2001,
-  HSA_EXT_STATUS_ERROR_INCOMPATIBLE_MODULE = 0x2002,
-  HSA_EXT_STATUS_ERROR_MODULE_ALREADY_INCLUDED = 0x2003,
-  HSA_EXT_STATUS_ERROR_SYMBOL_MISMATCH = 0x2004,
-  HSA_EXT_STATUS_ERROR_FINALIZATION_FAILED = 0x2005,
-  HSA_EXT_STATUS_ERROR_DIRECTIVE_MISMATCH = 0x2006
-};
-hsa_status_t hsa_ext_sampler_destroy (hsa_agent_t agent,
-				      hsa_ext_sampler_t sampler);
-hsa_status_t hsa_ext_program_finalize (
-  hsa_ext_program_t program, hsa_isa_t isa, int32_t call_convention,
-  hsa_ext_control_directives_t control_directives, const char *options,
-  hsa_code_object_type_t code_object_type, hsa_code_object_t *code_object);
-hsa_status_t hsa_ext_image_create (
-  hsa_agent_t agent, const hsa_ext_image_descriptor_t *image_descriptor,
-  const void *image_data, hsa_access_permission_t access_permission,
-  hsa_ext_image_t *image);
-hsa_status_t hsa_ext_program_destroy (hsa_ext_program_t program);
-hsa_status_t hsa_ext_image_copy (hsa_agent_t agent, hsa_ext_image_t src_image,
-				 const hsa_dim3_t *src_offset,
-				 hsa_ext_image_t dst_image,
-				 const hsa_dim3_t *dst_offset,
-				 const hsa_dim3_t *range);
-hsa_status_t hsa_ext_image_clear (hsa_agent_t agent, hsa_ext_image_t image,
-				  const void *data,
-				  const hsa_ext_image_region_t *image_region);
-enum
-{
-  HSA_EXT_AGENT_INFO_IMAGE_1D_MAX_ELEMENTS = 0x3000,
-  HSA_EXT_AGENT_INFO_IMAGE_1DA_MAX_ELEMENTS = 0x3001,
-  HSA_EXT_AGENT_INFO_IMAGE_1DB_MAX_ELEMENTS = 0x3002,
-  HSA_EXT_AGENT_INFO_IMAGE_2D_MAX_ELEMENTS = 0x3003,
-  HSA_EXT_AGENT_INFO_IMAGE_2DA_MAX_ELEMENTS = 0x3004,
-  HSA_EXT_AGENT_INFO_IMAGE_2DDEPTH_MAX_ELEMENTS = 0x3005,
-  HSA_EXT_AGENT_INFO_IMAGE_2DADEPTH_MAX_ELEMENTS = 0x3006,
-  HSA_EXT_AGENT_INFO_IMAGE_3D_MAX_ELEMENTS = 0x3007,
-  HSA_EXT_AGENT_INFO_IMAGE_ARRAY_MAX_LAYERS = 0x3008,
-  HSA_EXT_AGENT_INFO_MAX_IMAGE_RD_HANDLES = 0x3009,
-  HSA_EXT_AGENT_INFO_MAX_IMAGE_RORW_HANDLES = 0x300A,
-  HSA_EXT_AGENT_INFO_MAX_SAMPLER_HANDLERS = 0x300B
-};
-hsa_status_t
-hsa_ext_sampler_create (hsa_agent_t agent,
-			const hsa_ext_sampler_descriptor_t *sampler_descriptor,
-			hsa_ext_sampler_t *sampler);
-
-#endif /* _HSA_EXT_FINALIZE_H */
diff --git a/libgomp/plugin/plugin-hsa.c b/libgomp/plugin/plugin-hsa.c
deleted file mode 100644
index abd3bc6..0000000
--- a/libgomp/plugin/plugin-hsa.c
+++ /dev/null
@@ -1,1871 +0,0 @@
-/* Plugin for HSAIL execution.
-
-   Copyright (C) 2013-2020 Free Software Foundation, Inc.
-
-   Contributed by Martin Jambor <mjambor@suse.cz> and
-   Martin Liska <mliska@suse.cz>.
-
-   This file is part of the GNU Offloading and Multi Processing Library
-   (libgomp).
-
-   Libgomp is free software; you can redistribute it and/or modify it
-   under the terms of the GNU General Public License as published by
-   the Free Software Foundation; either version 3, or (at your option)
-   any later version.
-
-   Libgomp is distributed in the hope that it will be useful, but WITHOUT ANY
-   WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS
-   FOR A PARTICULAR PURPOSE.  See the GNU General Public License for
-   more details.
-
-   Under Section 7 of GPL version 3, you are granted additional
-   permissions described in the GCC Runtime Library Exception, version
-   3.1, as published by the Free Software Foundation.
-
-   You should have received a copy of the GNU General Public License and
-   a copy of the GCC Runtime Library Exception along with this program;
-   see the files COPYING3 and COPYING.RUNTIME respectively.  If not, see
-   <http://www.gnu.org/licenses/>.  */
-
-#include "config.h"
-#include <stdint.h>
-#include <stdio.h>
-#include <stdlib.h>
-#include <string.h>
-#include <pthread.h>
-#ifdef HAVE_INTTYPES_H
-#include <inttypes.h>
-#endif
-#include <stdbool.h>
-#include <hsa.h>
-#include <plugin/hsa_ext_finalize.h>
-#include <dlfcn.h>
-#include "libgomp-plugin.h"
-#include "gomp-constants.h"
-#include "secure_getenv.h"
-
-#ifdef HAVE_INTTYPES_H
-typedef uint64_t print_uint64_t;
-#else
-#define PRIu64 "lu"
-typedef unsigned long print_uint64_t; 
-#endif
-
-/* 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_agent_get_info_fn) (hsa_agent_t agent,
-					 hsa_agent_info_t attribute,
-					 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_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);
-  uint64_t (*hsa_queue_add_write_index_release_fn) (const hsa_queue_t *queue,
-						    uint64_t value);
-  uint64_t (*hsa_queue_load_read_index_acquire_fn) (const hsa_queue_t *queue);
-  void (*hsa_signal_store_relaxed_fn) (hsa_signal_t signal,
-				       hsa_signal_value_t value);
-  void (*hsa_signal_store_release_fn) (hsa_signal_t signal,
-				       hsa_signal_value_t value);
-  hsa_signal_value_t (*hsa_signal_wait_acquire_fn)
-    (hsa_signal_t signal, hsa_signal_condition_t condition,
-     hsa_signal_value_t compare_value, uint64_t timeout_hint,
-     hsa_wait_state_t wait_state_hint);
-  hsa_signal_value_t (*hsa_signal_load_acquire_fn) (hsa_signal_t signal);
-  hsa_status_t (*hsa_queue_destroy_fn) (hsa_queue_t *queue);
-
-  /* HSA finalizer.  */
-  hsa_status_t (*hsa_ext_program_add_module_fn) (hsa_ext_program_t program,
-						 hsa_ext_module_t module);
-  hsa_status_t (*hsa_ext_program_create_fn)
-    (hsa_machine_model_t machine_model, hsa_profile_t profile,
-     hsa_default_float_rounding_mode_t default_float_rounding_mode,
-     const char *options, hsa_ext_program_t *program);
-  hsa_status_t (*hsa_ext_program_destroy_fn) (hsa_ext_program_t program);
-  hsa_status_t (*hsa_ext_program_finalize_fn)
-    (hsa_ext_program_t program,hsa_isa_t isa,
-     int32_t call_convention, hsa_ext_control_directives_t control_directives,
-     const char *options, hsa_code_object_type_t code_object_type,
-     hsa_code_object_t *code_object);
-};
-
-/* HSA runtime functions that are initialized in init_hsa_context.  */
-
-static struct hsa_runtime_fn_info hsa_fns;
-
-/* Keep the following GOMP prefixed structures in sync with respective parts of
-   the compiler.  */
-
-/* Structure describing the run-time and grid properties of an HSA kernel
-   lauch.  */
-
-struct GOMP_kernel_launch_attributes
-{
-  /* Number of dimensions the workload has.  Maximum number is 3.  */
-  uint32_t ndim;
-  /* Size of the grid in the three respective dimensions.  */
-  uint32_t gdims[3];
-  /* Size of work-groups in the respective dimensions.  */
-  uint32_t wdims[3];
-};
-
-/* Collection of information needed for a dispatch of a kernel from a
-   kernel.  */
-
-struct GOMP_hsa_kernel_dispatch
-{
-  /* Pointer to a command queue associated with a kernel dispatch agent.  */
-  void *queue;
-  /* Pointer to reserved memory for OMP data struct copying.  */
-  void *omp_data_memory;
-  /* Pointer to a memory space used for kernel arguments passing.  */
-  void *kernarg_address;
-  /* Kernel object.  */
-  uint64_t object;
-  /* Synchronization signal used for dispatch synchronization.  */
-  uint64_t signal;
-  /* Private segment size.  */
-  uint32_t private_segment_size;
-  /* Group segment size.  */
-  uint32_t group_segment_size;
-  /* Number of children kernel dispatches.  */
-  uint64_t kernel_dispatch_count;
-  /* Debug purpose argument.  */
-  uint64_t debug;
-  /* Levels-var ICV.  */
-  uint64_t omp_level;
-  /* Kernel dispatch structures created for children kernel dispatches.  */
-  struct GOMP_hsa_kernel_dispatch **children_dispatches;
-  /* Number of threads.  */
-  uint32_t omp_num_threads;
-};
-
-/* Part of the libgomp plugin interface.  Return the name of the accelerator,
-   which is "hsa".  */
-
-const char *
-GOMP_OFFLOAD_get_name (void)
-{
-  return "hsa";
-}
-
-/* Part of the libgomp plugin interface.  Return the specific capabilities the
-   HSA accelerator have.  */
-
-unsigned int
-GOMP_OFFLOAD_get_caps (void)
-{
-  return GOMP_OFFLOAD_CAP_SHARED_MEM | GOMP_OFFLOAD_CAP_OPENMP_400;
-}
-
-/* Part of the libgomp plugin interface.  Identify as HSA accelerator.  */
-
-int
-GOMP_OFFLOAD_get_type (void)
-{
-  return OFFLOAD_TARGET_TYPE_HSA;
-}
-
-/* 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;
-}
-
-/* Flag to decide whether print to stderr information about what is going on.
-   Set in init_debug depending on environment variables.  */
-
-static bool debug;
-
-/* Flag to decide if the runtime should suppress a possible fallback to host
-   execution.  */
-
-static bool suppress_host_fallback;
-
-/* Flag to locate HSA runtime shared library that is dlopened
-   by this plug-in.  */
-
-static const char *hsa_runtime_lib;
-
-/* Flag to decide if the runtime should support also CPU devices (can be
-   a simulator).  */
-
-static bool support_cpu_devices;
-
-/* Initialize debug and suppress_host_fallback according to the environment.  */
-
-static void
-init_enviroment_variables (void)
-{
-  if (secure_getenv ("HSA_DEBUG"))
-    debug = true;
-  else
-    debug = false;
-
-  if (secure_getenv ("HSA_SUPPRESS_HOST_FALLBACK"))
-    suppress_host_fallback = true;
-  else
-    suppress_host_fallback = false;
-
-  hsa_runtime_lib = secure_getenv ("HSA_RUNTIME_LIB");
-  if (hsa_runtime_lib == NULL)
-    hsa_runtime_lib = HSA_RUNTIME_LIB "libhsa-runtime64.so";
-
-  support_cpu_devices = secure_getenv ("HSA_SUPPORT_CPU_DEVICES");
-}
-
-/* Print a logging message with PREFIX to stderr if HSA_DEBUG value
-   is set to true.  */
-
-#define HSA_LOG(prefix, ...) \
-  do \
-  { \
-    if (debug) \
-      { \
-	fprintf (stderr, prefix); \
-	fprintf (stderr, __VA_ARGS__); \
-      } \
-  } \
-  while (false)
-
-/* Print a debugging message to stderr.  */
-
-#define HSA_DEBUG(...) HSA_LOG ("HSA debug: ", __VA_ARGS__)
-
-/* Print a warning message to stderr.  */
-
-#define HSA_WARNING(...) HSA_LOG ("HSA 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, "HSA warning: %s\nRuntime message: %s", 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 ("HSA fatal error: %s\nRuntime message: %s", 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 ("HSA fatal error: %s\nRuntime message: %s", str,
-		     hsa_error_msg);
-  return false;
-}
-
-struct hsa_kernel_description
-{
-  const char *name;
-  unsigned omp_data_size;
-  bool gridified_kernel_p;
-  unsigned kernel_dependencies_count;
-  const char **kernel_dependencies;
-};
-
-struct global_var_info
-{
-  const char *name;
-  void *address;
-};
-
-/* Data passed by the static initializer of a compilation unit containing BRIG
-   to GOMP_offload_register.  */
-
-struct brig_image_desc
-{
-  hsa_ext_module_t brig_module;
-  const unsigned kernel_count;
-  struct hsa_kernel_description *kernel_infos;
-  const unsigned global_variable_count;
-  struct global_var_info *global_variables;
-};
-
-struct agent_info;
-
-/* Information required to identify, finalize and run any given kernel.  */
-
-struct kernel_info
-{
-  /* Name of the kernel, required to locate it within the brig module.  */
-  const char *name;
-  /* Size of memory space for OMP data.  */
-  unsigned omp_data_size;
-  /* The specific agent the kernel has been or will be finalized for and run
-     on.  */
-  struct agent_info *agent;
-  /* The specific module where the kernel takes place.  */
-  struct module_info *module;
-  /* Mutex enforcing that at most once thread ever initializes a kernel for
-     use.  A thread should have locked agent->modules_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;
-  /* List of all kernel dependencies.  */
-  const char **dependencies;
-  /* Number of dependencies.  */
-  unsigned dependencies_count;
-  /* Maximum OMP data size necessary for kernel from kernel dispatches.  */
-  unsigned max_omp_data_size;
-  /* True if the kernel is gridified.  */
-  bool gridified_kernel_p;
-};
-
-/* Information about a particular brig module, its image and kernels.  */
-
-struct module_info
-{
-  /* The next and previous module in the linked list of modules of an agent.  */
-  struct module_info *next, *prev;
-  /* The description with which the program has registered the image.  */
-  struct brig_image_desc *image_desc;
-
-  /* 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[];
-};
-
-/* Information about shared brig library.  */
-
-struct brig_library_info
-{
-  char *file_name;
-  hsa_ext_module_t image;
-};
-
-/* Description of an HSA GPU agent 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;
-  /* Whether the agent has been initialized.  The fields below are usable only
-     if it has been.  */
-  bool initialized;
-  /* The HSA ISA of this agent.  */
-  hsa_isa_t isa;
-  /* Command queue of the agent.  */
-  hsa_queue_t *command_q;
-  /* Kernel from kernel dispatch command queue.  */
-  hsa_queue_t *kernel_dispatch_command_q;
-  /* The HSA memory region from which to allocate kernel arguments.  */
-  hsa_region_t kernarg_region;
-
-  /* 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 modules_rwlock;
-  /* The first module in a linked list of modules associated with this
-     kernel.  */
-  struct module_info *first_module;
-
-  /* Mutex enforcing that only one thread will finalize the HSA program.  A
-     thread should have locked agent->modules_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;
-  /* Flag whether the program was finalized but with a failure.  */
-  bool prog_finalized_error;
-  /* HSA executable - the finalized program that is used to locate kernels.  */
-  hsa_executable_t executable;
-  /* List of BRIG libraries.  */
-  struct brig_library_info **brig_libraries;
-  /* Number of loaded shared BRIG libraries.  */
-  unsigned brig_libraries_count;
-};
-
-/* 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;
-};
-
-/* Information about the whole HSA environment and all of its agents.  */
-
-static struct hsa_context_info hsa_context;
-
-#define DLSYM_FN(function) \
-  hsa_fns.function##_fn = dlsym (handle, #function); \
-  if (hsa_fns.function##_fn == NULL) \
-    goto dl_fail;
-
-static bool
-init_hsa_runtime_functions (void)
-{
-  void *handle = dlopen (hsa_runtime_lib, RTLD_LAZY);
-  if (handle == NULL)
-    goto dl_fail;
-
-  DLSYM_FN (hsa_status_string)
-  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_free)
-  DLSYM_FN (hsa_signal_destroy)
-  DLSYM_FN (hsa_executable_get_symbol)
-  DLSYM_FN (hsa_executable_symbol_get_info)
-  DLSYM_FN (hsa_queue_add_write_index_release)
-  DLSYM_FN (hsa_queue_load_read_index_acquire)
-  DLSYM_FN (hsa_signal_wait_acquire)
-  DLSYM_FN (hsa_signal_store_relaxed)
-  DLSYM_FN (hsa_signal_store_release)
-  DLSYM_FN (hsa_signal_load_acquire)
-  DLSYM_FN (hsa_queue_destroy)
-  DLSYM_FN (hsa_ext_program_add_module)
-  DLSYM_FN (hsa_ext_program_create)
-  DLSYM_FN (hsa_ext_program_destroy)
-  DLSYM_FN (hsa_ext_program_finalize)
-  return true;
-
- dl_fail:
-  HSA_DEBUG ("while loading %s: %s\n", hsa_runtime_lib, dlerror ());
-  return false;
-}
-
-/* Find kernel for an AGENT by name provided in KERNEL_NAME.  */
-
-static struct kernel_info *
-get_kernel_for_agent (struct agent_info *agent, const char *kernel_name)
-{
-  struct module_info *module = agent->first_module;
-
-  while (module)
-    {
-      for (unsigned i = 0; i < module->kernel_count; i++)
-	if (strcmp (module->kernels[i].name, kernel_name) == 0)
-	  return &module->kernels[i];
-
-      module = module->next;
-    }
-
-  return NULL;
-}
-
-/* Return true if the agent is a GPU and acceptable of concurrent submissions
-   from different threads.  */
-
-static bool
-suitable_hsa_agent_p (hsa_agent_t agent)
-{
-  hsa_device_type_t device_type;
-  hsa_status_t status
-    = hsa_fns.hsa_agent_get_info_fn (agent, HSA_AGENT_INFO_DEVICE,
-				     &device_type);
-  if (status != HSA_STATUS_SUCCESS)
-    return false;
-
-  switch (device_type)
-    {
-    case HSA_DEVICE_TYPE_GPU:
-      break;
-    case HSA_DEVICE_TYPE_CPU:
-      if (!support_cpu_devices)
-	return false;
-      break;
-    default:
-      return false;
-    }
-
-  uint32_t features = 0;
-  status = hsa_fns.hsa_agent_get_info_fn (agent, HSA_AGENT_INFO_FEATURE,
-					  &features);
-  if (status != HSA_STATUS_SUCCESS
-      || !(features & HSA_AGENT_FEATURE_KERNEL_DISPATCH))
-    return false;
-  hsa_queue_type_t queue_type;
-  status = hsa_fns.hsa_agent_get_info_fn (agent, HSA_AGENT_INFO_QUEUE_TYPE,
-					  &queue_type);
-  if (status != HSA_STATUS_SUCCESS
-      || (queue_type != HSA_QUEUE_TYPE_MULTI))
-    return false;
-
-  return true;
-}
-
-/* Callback of hsa_iterate_agents, if AGENT is a GPU device, increment
-   agent_count in hsa_context.  */
-
-static hsa_status_t
-count_gpu_agents (hsa_agent_t agent, void *data __attribute__ ((unused)))
-{
-  if (suitable_hsa_agent_p (agent))
-    hsa_context.agent_count++;
-  return HSA_STATUS_SUCCESS;
-}
-
-/* Callback of hsa_iterate_agents, if AGENT is a GPU device, assign the agent
-   id to the describing structure in the hsa context.  The index of the
-   structure is pointed to by DATA, increment it afterwards.  */
-
-static hsa_status_t
-assign_agent_ids (hsa_agent_t agent, void *data)
-{
-  if (suitable_hsa_agent_p (agent))
-    {
-      int *agent_index = (int *) data;
-      hsa_context.agents[*agent_index].id = agent;
-      ++*agent_index;
-    }
-  return HSA_STATUS_SUCCESS;
-}
-
-/* Initialize hsa_context if it has not already been done.
-   Return TRUE on success.  */
-
-static bool
-init_hsa_context (void)
-{
-  hsa_status_t status;
-  int agent_index = 0;
-
-  if (hsa_context.initialized)
-    return true;
-  init_enviroment_variables ();
-  if (!init_hsa_runtime_functions ())
-    {
-      HSA_DEBUG ("Run-time could not be dynamically opened\n");
-      return false;
-    }
-  status = hsa_fns.hsa_init_fn ();
-  if (status != HSA_STATUS_SUCCESS)
-    return hsa_error ("Run-time could not be initialized", status);
-  HSA_DEBUG ("HSA run-time initialized\n");
-  status = hsa_fns.hsa_iterate_agents_fn (count_gpu_agents, NULL);
-  if (status != HSA_STATUS_SUCCESS)
-    return hsa_error ("HSA GPU devices could not be enumerated", status);
-  HSA_DEBUG ("There are %i HSA GPU devices.\n", hsa_context.agent_count);
-
-  hsa_context.agents
-    = GOMP_PLUGIN_malloc_cleared (hsa_context.agent_count
-				  * sizeof (struct agent_info));
-  status = hsa_fns.hsa_iterate_agents_fn (assign_agent_ids, &agent_index);
-  if (agent_index != hsa_context.agent_count)
-    {
-      GOMP_PLUGIN_error ("Failed to assign IDs to all HSA agents");
-      return false;
-    }
-  hsa_context.initialized = true;
-  return true;
-}
-
-/* Callback of dispatch queues to report errors.  */
-
-static void
-queue_callback (hsa_status_t status,
-		hsa_queue_t *queue __attribute__ ((unused)),
-		void *data __attribute__ ((unused)))
-{
-  hsa_fatal ("Asynchronous queue error", status);
-}
-
-/* Callback of hsa_agent_iterate_regions.  Determine if a memory REGION can be
-   used for kernarg allocations and if so write it to the memory pointed to by
-   DATA and break the query.  */
-
-static hsa_status_t
-get_kernarg_memory_region (hsa_region_t region, void *data)
-{
-  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 & HSA_REGION_GLOBAL_FLAG_KERNARG)
-    {
-      hsa_region_t *ret = (hsa_region_t *) data;
-      *ret = region;
-      return HSA_STATUS_INFO_BREAK;
-    }
-  return HSA_STATUS_SUCCESS;
-}
-
-/* Part of the libgomp plugin interface.  Return the number of HSA devices on
-   the system.  */
-
-int
-GOMP_OFFLOAD_get_num_devices (void)
-{
-  if (!init_hsa_context ())
-    return 0;
-  return hsa_context.agent_count;
-}
-
-/* Part of the libgomp plugin interface.  Initialize agent number N so that it
-   can be used for computation.  Return TRUE on success.  */
-
-bool
-GOMP_OFFLOAD_init_device (int n)
-{
-  if (!init_hsa_context ())
-    return false;
-  if (n >= hsa_context.agent_count)
-    {
-      GOMP_PLUGIN_error ("Request to initialize non-existing HSA device %i", n);
-      return false;
-    }
-  struct agent_info *agent = &hsa_context.agents[n];
-
-  if (agent->initialized)
-    return true;
-
-  if (pthread_rwlock_init (&agent->modules_rwlock, NULL))
-    {
-      GOMP_PLUGIN_error ("Failed to initialize an HSA agent rwlock");
-      return false;
-    }
-  if (pthread_mutex_init (&agent->prog_mutex, NULL))
-    {
-      GOMP_PLUGIN_error ("Failed to initialize an HSA agent program mutex");
-      return false;
-    }
-
-  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 HSA agent",
-    	   	      status);
-  status = hsa_fns.hsa_agent_get_info_fn (agent->id, HSA_AGENT_INFO_ISA,
-					  &agent->isa);
-  if (status != HSA_STATUS_SUCCESS)
-    return hsa_error ("Error querying the ISA of the agent", status);
-  status = hsa_fns.hsa_queue_create_fn (agent->id, queue_size,
-					HSA_QUEUE_TYPE_MULTI,
-					queue_callback, NULL, UINT32_MAX,
-					UINT32_MAX,
-					&agent->command_q);
-  if (status != HSA_STATUS_SUCCESS)
-    return hsa_error ("Error creating command queue", status);
-
-  status = hsa_fns.hsa_queue_create_fn (agent->id, queue_size,
-					HSA_QUEUE_TYPE_MULTI,
-					queue_callback, NULL, UINT32_MAX,
-					UINT32_MAX,
-					&agent->kernel_dispatch_command_q);
-  if (status != HSA_STATUS_SUCCESS)
-    return hsa_error ("Error creating kernel dispatch command queue", status);
-
-  agent->kernarg_region.handle = (uint64_t) -1;
-  status = hsa_fns.hsa_agent_iterate_regions_fn (agent->id,
-						 get_kernarg_memory_region,
-						 &agent->kernarg_region);
-  if (agent->kernarg_region.handle == (uint64_t) -1)
-    {
-      GOMP_PLUGIN_error ("Could not find suitable memory region for kernel "
-			 "arguments");
-      return false;
-    }
-  HSA_DEBUG ("HSA agent initialized, queue has id %llu\n",
-	     (long long unsigned) agent->command_q->id);
-  HSA_DEBUG ("HSA agent initialized, kernel dispatch queue has id %llu\n",
-	     (long long unsigned) agent->kernel_dispatch_command_q->id);
-  agent->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 HSA context.");
-      return NULL;
-    }
-  if (n >= hsa_context.agent_count)
-    {
-      GOMP_PLUGIN_error ("Request to operate on anon-existing HSA device %i", n);
-      return NULL;
-    }
-  if (!hsa_context.agents[n].initialized)
-    {
-      GOMP_PLUGIN_error ("Attempt to use an uninitialized HSA agent.");
-      return NULL;
-    }
-  return &hsa_context.agents[n];
-}
-
-/* Insert MODULE to the linked list of modules of AGENT.  */
-
-static void
-add_module_to_agent (struct agent_info *agent, struct module_info *module)
-{
-  if (agent->first_module)
-    agent->first_module->prev = module;
-  module->next = agent->first_module;
-  module->prev = NULL;
-  agent->first_module = module;
-}
-
-/* Remove MODULE from the linked list of modules of AGENT.  */
-
-static void
-remove_module_from_agent (struct agent_info *agent, struct module_info *module)
-{
-  if (agent->first_module == module)
-    agent->first_module = module->next;
-  if (module->prev)
-    module->prev->next = module->next;
-  if (module->next)
-    module->next->prev = module->prev;
-}
-
-/* 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 || agent->prog_finalized_error)
-    return true;
-
-  hsa_status_t status;
-
-  HSA_DEBUG ("Destroying the current HSA program.\n");
-
-  status = hsa_fns.hsa_executable_destroy_fn (agent->executable);
-  if (status != HSA_STATUS_SUCCESS)
-    return hsa_error ("Could not destroy HSA executable", status);
-
-  struct module_info *module;
-  for (module = agent->first_module; module; module = module->next)
-    {
-      int i;
-      for (i = 0; i < module->kernel_count; i++)
-	module->kernels[i].initialized = false;
-    }
-  agent->prog_finalized = false;
-  return true;
-}
-
-/* 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->omp_data_size = d->omp_data_size;
-  kernel->gridified_kernel_p = d->gridified_kernel_p;
-  kernel->dependencies_count = d->kernel_dependencies_count;
-  kernel->dependencies = d->kernel_dependencies;
-  if (pthread_mutex_init (&kernel->init_mutex, NULL))
-    {
-      GOMP_PLUGIN_error ("Failed to initialize an HSA kernel mutex");
-      return false;
-    }
-  return true;
-}
-
-/* Part of the libgomp plugin interface.  Load BRIG module described by struct
-   brig_image_desc in TARGET_DATA and return references to kernel descriptors
-   in TARGET_TABLE.  */
-
-int
-GOMP_OFFLOAD_load_image (int ord, unsigned version, const void *target_data,
-			 struct addr_pair **target_table)
-{
-  if (GOMP_VERSION_DEV (version) > GOMP_VERSION_HSA)
-    {
-      GOMP_PLUGIN_error ("Offload data incompatible with HSA plugin"
-			 " (expected %u, received %u)",
-			 GOMP_VERSION_HSA, GOMP_VERSION_DEV (version));
-      return -1;
-    }
-
-  struct brig_image_desc *image_desc = (struct brig_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;
-
-  agent = get_agent_info (ord);
-  if (!agent)
-    return -1;
-
-  if (pthread_rwlock_wrlock (&agent->modules_rwlock))
-    {
-      GOMP_PLUGIN_error ("Unable to write-lock an HSA agent rwlock");
-      return -1;
-    }
-  if (agent->prog_finalized
-      && !destroy_hsa_program (agent))
-    return -1;
-
-  HSA_DEBUG ("Encountered %d kernels in an image\n", kernel_count);
-  pair = GOMP_PLUGIN_malloc (kernel_count * 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;
-
-  kernel = &module->kernels[0];
-
-  /* Allocate memory for kernel dependencies.  */
-  for (unsigned i = 0; i < kernel_count; i++)
-    {
-      pair->start = (uintptr_t) kernel;
-      pair->end = (uintptr_t) (kernel + 1);
-
-      struct hsa_kernel_description *d = &image_desc->kernel_infos[i];
-      if (!init_basic_kernel_info (kernel, d, agent, module))
-	return -1;
-      kernel++;
-      pair++;
-    }
-
-  add_module_to_agent (agent, module);
-  if (pthread_rwlock_unlock (&agent->modules_rwlock))
-    {
-      GOMP_PLUGIN_error ("Unable to unlock an HSA agent rwlock");
-      return -1;
-    }
-  return kernel_count;
-}
-
-/* Add a shared BRIG library from a FILE_NAME to an AGENT.  */
-
-static struct brig_library_info *
-add_shared_library (const char *file_name, struct agent_info *agent)
-{
-  struct brig_library_info *library = NULL;
-
-  void *f = dlopen (file_name, RTLD_NOW);
-  void *start = dlsym (f, "__brig_start");
-  void *end = dlsym (f, "__brig_end");
-
-  if (start == NULL || end == NULL)
-    return NULL;
-
-  unsigned size = end - start;
-  char *buf = (char *) GOMP_PLUGIN_malloc (size);
-  memcpy (buf, start, size);
-
-  library = GOMP_PLUGIN_malloc (sizeof (struct agent_info));
-  library->file_name = (char *) GOMP_PLUGIN_malloc
-    ((strlen (file_name) + 1));
-  strcpy (library->file_name, file_name);
-  library->image = (hsa_ext_module_t) buf;
-
-  return library;
-}
-
-/* Release memory used for BRIG shared libraries that correspond
-   to an AGENT.  */
-
-static void
-release_agent_shared_libraries (struct agent_info *agent)
-{
-  for (unsigned i = 0; i < agent->brig_libraries_count; i++)
-    if (agent->brig_libraries[i])
-      {
-	free (agent->brig_libraries[i]->file_name);
-	free (agent->brig_libraries[i]->image);
-	free (agent->brig_libraries[i]);
-      }
-
-  free (agent->brig_libraries);
-}
-
-/* Create and finalize the program consisting of all loaded modules.  */
-
-static void
-create_and_finalize_hsa_program (struct agent_info *agent)
-{
-  hsa_status_t status;
-  hsa_ext_program_t prog_handle;
-  int mi = 0;
-
-  if (pthread_mutex_lock (&agent->prog_mutex))
-    GOMP_PLUGIN_fatal ("Could not lock an HSA agent program mutex");
-  if (agent->prog_finalized)
-    goto final;
-
-  status = hsa_fns.hsa_ext_program_create_fn
-    (HSA_MACHINE_MODEL_LARGE, HSA_PROFILE_FULL,
-     HSA_DEFAULT_FLOAT_ROUNDING_MODE_DEFAULT,
-     NULL, &prog_handle);
-  if (status != HSA_STATUS_SUCCESS)
-    hsa_fatal ("Could not create an HSA program", status);
-
-  HSA_DEBUG ("Created a finalized program\n");
-
-  struct module_info *module = agent->first_module;
-  while (module)
-    {
-      status = hsa_fns.hsa_ext_program_add_module_fn
-	(prog_handle, module->image_desc->brig_module);
-      if (status != HSA_STATUS_SUCCESS)
-	hsa_fatal ("Could not add a module to the HSA program", status);
-      module = module->next;
-      mi++;
-    }
-
-  /* Load all shared libraries.  */
-  const char *libraries[] = { "libhsamath.so", "libhsastd.so" };
-  const unsigned libraries_count = sizeof (libraries) / sizeof (const char *);
-
-  agent->brig_libraries_count = libraries_count;
-  agent->brig_libraries = GOMP_PLUGIN_malloc_cleared
-    (sizeof (struct brig_library_info) * libraries_count);
-
-  for (unsigned i = 0; i < libraries_count; i++)
-    {
-      struct brig_library_info *library = add_shared_library (libraries[i],
-							      agent);
-      if (library == NULL)
-	{
-	  HSA_WARNING ("Could not open a shared BRIG library: %s\n",
-		       libraries[i]);
-	  continue;
-	}
-
-      status = hsa_fns.hsa_ext_program_add_module_fn (prog_handle,
-						      library->image);
-      if (status != HSA_STATUS_SUCCESS)
-	hsa_warn ("Could not add a shared BRIG library the HSA program",
-		  status);
-      else
-	HSA_DEBUG ("a shared BRIG library has been added to a program: %s\n",
-		   libraries[i]);
-    }
-
-  hsa_ext_control_directives_t control_directives;
-  memset (&control_directives, 0, sizeof (control_directives));
-  hsa_code_object_t code_object;
-  status = hsa_fns.hsa_ext_program_finalize_fn
-    (prog_handle, agent->isa,HSA_EXT_FINALIZER_CALL_CONVENTION_AUTO,
-     control_directives, "", HSA_CODE_OBJECT_TYPE_PROGRAM, &code_object);
-  if (status != HSA_STATUS_SUCCESS)
-    {
-      hsa_warn ("Finalization of the HSA program failed", status);
-      goto failure;
-    }
-
-  HSA_DEBUG ("Finalization done\n");
-  hsa_fns.hsa_ext_program_destroy_fn (prog_handle);
-
-  status
-    = hsa_fns.hsa_executable_create_fn (HSA_PROFILE_FULL,
-					HSA_EXECUTABLE_STATE_UNFROZEN,
-					"", &agent->executable);
-  if (status != HSA_STATUS_SUCCESS)
-    hsa_fatal ("Could not create HSA executable", status);
-
-  module = agent->first_module;
-  while (module)
-    {
-      /* Initialize all global variables declared in the module.  */
-      for (unsigned i = 0; i < module->image_desc->global_variable_count; i++)
-	{
-	  struct global_var_info *var;
-	  var = &module->image_desc->global_variables[i];
-	  status = hsa_fns.hsa_executable_global_variable_define_fn
-	    (agent->executable, var->name, var->address);
-
-	  HSA_DEBUG ("Defining global variable: %s, address: %p\n", var->name,
-		     var->address);
-
-	  if (status != HSA_STATUS_SUCCESS)
-	    hsa_fatal ("Could not define a global variable in the HSA program",
-		       status);
-	}
-
-      module = module->next;
-    }
-
-  status = hsa_fns.hsa_executable_load_code_object_fn (agent->executable,
-						       agent->id,
-						       code_object, "");
-  if (status != HSA_STATUS_SUCCESS)
-    hsa_fatal ("Could not add a code object to the HSA executable", status);
-  status = hsa_fns.hsa_executable_freeze_fn (agent->executable, "");
-  if (status != HSA_STATUS_SUCCESS)
-    hsa_fatal ("Could not freeze the HSA executable", status);
-
-  HSA_DEBUG ("Froze HSA executable with the finalized code object\n");
-
-  /* If all goes good, jump to final.  */
-  goto final;
-
-failure:
-  agent->prog_finalized_error = true;
-
-final:
-  agent->prog_finalized = true;
-
-  if (pthread_mutex_unlock (&agent->prog_mutex))
-    GOMP_PLUGIN_fatal ("Could not unlock an HSA agent program mutex");
-}
-
-/* Create kernel dispatch data structure for given KERNEL.  */
-
-static struct GOMP_hsa_kernel_dispatch *
-create_single_kernel_dispatch (struct kernel_info *kernel,
-			       unsigned omp_data_size)
-{
-  struct agent_info *agent = kernel->agent;
-  struct GOMP_hsa_kernel_dispatch *shadow
-    = GOMP_PLUGIN_malloc_cleared (sizeof (struct GOMP_hsa_kernel_dispatch));
-
-  shadow->queue = agent->command_q;
-  shadow->omp_data_memory
-    = omp_data_size > 0 ? GOMP_PLUGIN_malloc (omp_data_size) : NULL;
-  unsigned dispatch_count = kernel->dependencies_count;
-  shadow->kernel_dispatch_count = dispatch_count;
-
-  shadow->children_dispatches
-    = GOMP_PLUGIN_malloc (dispatch_count * sizeof (shadow));
-
-  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 HSA sync signal", status);
-
-  shadow->signal = sync_signal.handle;
-  shadow->private_segment_size = kernel->private_segment_size;
-  shadow->group_segment_size = kernel->group_segment_size;
-
-  status
-    = hsa_fns.hsa_memory_allocate_fn (agent->kernarg_region,
-				      kernel->kernarg_segment_size,
-				      &shadow->kernarg_address);
-  if (status != HSA_STATUS_SUCCESS)
-    hsa_fatal ("Could not allocate memory for HSA kernel arguments", status);
-
-  return shadow;
-}
-
-/* Release data structure created for a kernel dispatch in SHADOW argument.  */
-
-static void
-release_kernel_dispatch (struct GOMP_hsa_kernel_dispatch *shadow)
-{
-  HSA_DEBUG ("Released kernel dispatch: %p has value: %" PRIu64 " (%p)\n",
-	     shadow, (print_uint64_t) shadow->debug,
-	     (void *) (uintptr_t) shadow->debug);
-
-  hsa_fns.hsa_memory_free_fn (shadow->kernarg_address);
-
-  hsa_signal_t s;
-  s.handle = shadow->signal;
-  hsa_fns.hsa_signal_destroy_fn (s);
-
-  free (shadow->omp_data_memory);
-
-  for (unsigned i = 0; i < shadow->kernel_dispatch_count; i++)
-    release_kernel_dispatch (shadow->children_dispatches[i]);
-
-  free (shadow->children_dispatches);
-  free (shadow);
-}
-
-/* Initialize a KERNEL without its dependencies.  MAX_OMP_DATA_SIZE is used
-   to calculate maximum necessary memory for OMP data allocation.  */
-
-static void
-init_single_kernel (struct kernel_info *kernel, unsigned *max_omp_data_size)
-{
-  hsa_status_t status;
-  struct agent_info *agent = kernel->agent;
-  hsa_executable_symbol_t kernel_symbol;
-  status = hsa_fns.hsa_executable_get_symbol_fn (agent->executable, NULL,
-						 kernel->name, agent->id,
-						 0, &kernel_symbol);
-  if (status != HSA_STATUS_SUCCESS)
-    {
-      hsa_warn ("Could not find symbol for kernel in the code object", status);
-      goto failure;
-    }
-  HSA_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);
-
-  HSA_DEBUG ("Kernel structure for %s fully initialized with "
-	     "following segment sizes: \n", kernel->name);
-  HSA_DEBUG ("  group_segment_size: %u\n",
-	     (unsigned) kernel->group_segment_size);
-  HSA_DEBUG ("  private_segment_size: %u\n",
-	     (unsigned) kernel->private_segment_size);
-  HSA_DEBUG ("  kernarg_segment_size: %u\n",
-	     (unsigned) kernel->kernarg_segment_size);
-  HSA_DEBUG ("  omp_data_size: %u\n", kernel->omp_data_size);
-  HSA_DEBUG ("  gridified_kernel_p: %u\n", kernel->gridified_kernel_p);
-
-  if (kernel->omp_data_size > *max_omp_data_size)
-    *max_omp_data_size = kernel->omp_data_size;
-
-  for (unsigned i = 0; i < kernel->dependencies_count; i++)
-    {
-      struct kernel_info *dependency
-	= get_kernel_for_agent (agent, kernel->dependencies[i]);
-
-      if (dependency == NULL)
-	{
-	  HSA_DEBUG ("Could not find a dependency for a kernel: %s, "
-		     "dependency name: %s\n", kernel->name,
-		     kernel->dependencies[i]);
-	  goto failure;
-	}
-
-      if (dependency->dependencies_count > 0)
-	{
-	  HSA_DEBUG ("HSA does not allow kernel dispatching code with "
-		     "a depth bigger than one\n");
-	  goto failure;
-	}
-
-      init_single_kernel (dependency, max_omp_data_size);
-    }
-
-  return;
-
-failure:
-  kernel->initialization_failed = true;
-}
-
-/* Indent stream F by INDENT spaces.  */
-
-static void
-indent_stream (FILE *f, unsigned indent)
-{
-  fprintf (f, "%*s", indent, "");
-}
-
-/* Dump kernel DISPATCH data structure and indent it by INDENT spaces.  */
-
-static void
-print_kernel_dispatch (struct GOMP_hsa_kernel_dispatch *dispatch, unsigned indent)
-{
-  indent_stream (stderr, indent);
-  fprintf (stderr, "this: %p\n", dispatch);
-  indent_stream (stderr, indent);
-  fprintf (stderr, "queue: %p\n", dispatch->queue);
-  indent_stream (stderr, indent);
-  fprintf (stderr, "omp_data_memory: %p\n", dispatch->omp_data_memory);
-  indent_stream (stderr, indent);
-  fprintf (stderr, "kernarg_address: %p\n", dispatch->kernarg_address);
-  indent_stream (stderr, indent);
-  fprintf (stderr, "object: %" PRIu64 "\n", (print_uint64_t) dispatch->object);
-  indent_stream (stderr, indent);
-  fprintf (stderr, "signal: %" PRIu64 "\n", (print_uint64_t) dispatch->signal);
-  indent_stream (stderr, indent);
-  fprintf (stderr, "private_segment_size: %u\n",
-	   dispatch->private_segment_size);
-  indent_stream (stderr, indent);
-  fprintf (stderr, "group_segment_size: %u\n",
-	   dispatch->group_segment_size);
-  indent_stream (stderr, indent);
-  fprintf (stderr, "children dispatches: %" PRIu64 "\n",
-	   (print_uint64_t) dispatch->kernel_dispatch_count);
-  indent_stream (stderr, indent);
-  fprintf (stderr, "omp_num_threads: %u\n",
-	   dispatch->omp_num_threads);
-  fprintf (stderr, "\n");
-
-  for (unsigned i = 0; i < dispatch->kernel_dispatch_count; i++)
-    print_kernel_dispatch (dispatch->children_dispatches[i], indent + 2);
-}
-
-/* Create kernel dispatch data structure for a KERNEL and all its
-   dependencies.  */
-
-static struct GOMP_hsa_kernel_dispatch *
-create_kernel_dispatch (struct kernel_info *kernel, unsigned omp_data_size)
-{
-  struct GOMP_hsa_kernel_dispatch *shadow
-    = create_single_kernel_dispatch (kernel, omp_data_size);
-  shadow->omp_num_threads = 64;
-  shadow->debug = 0;
-  shadow->omp_level = kernel->gridified_kernel_p ? 1 : 0;
-
-  /* Create kernel dispatch data structures.  We do not allow to have
-     a kernel dispatch with depth bigger than one.  */
-  for (unsigned i = 0; i < kernel->dependencies_count; i++)
-    {
-      struct kernel_info *dependency
-	= get_kernel_for_agent (kernel->agent, kernel->dependencies[i]);
-      shadow->children_dispatches[i]
-	= create_single_kernel_dispatch (dependency, omp_data_size);
-      shadow->children_dispatches[i]->queue
-	= kernel->agent->kernel_dispatch_command_q;
-      shadow->children_dispatches[i]->omp_level = 1;
-    }
-
-  return shadow;
-}
-
-/* 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 an HSA kernel initialization mutex");
-  if (kernel->initialized)
-    {
-      if (pthread_mutex_unlock (&kernel->init_mutex))
-	GOMP_PLUGIN_fatal ("Could not unlock an HSA kernel initialization "
-			   "mutex");
-
-      return;
-    }
-
-  /* Precomputed maximum size of OMP data necessary for a kernel from kernel
-     dispatch operation.  */
-  init_single_kernel (kernel, &kernel->max_omp_data_size);
-
-  if (!kernel->initialization_failed)
-    HSA_DEBUG ("\n");
-
-  kernel->initialized = true;
-  if (pthread_mutex_unlock (&kernel->init_mutex))
-    GOMP_PLUGIN_fatal ("Could not unlock an HSA kernel initialization "
-		       "mutex");
-}
-
-/* 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.  */
-
-static bool
-parse_target_attributes (void **input,
-			 struct GOMP_kernel_launch_attributes *def,
-			 struct GOMP_kernel_launch_attributes **result)
-{
-  if (!input)
-    GOMP_PLUGIN_fatal ("No target arguments provided");
-
-  bool attrs_found = false;
-  while (*input)
-    {
-      uintptr_t id = (uintptr_t) *input;
-      if ((id & GOMP_TARGET_ARG_DEVICE_MASK) == GOMP_DEVICE_HSA
-	  && ((id & GOMP_TARGET_ARG_ID_MASK)
-	      == GOMP_TARGET_ARG_HSA_KERNEL_ATTRIBUTES))
-	{
-	  input++;
-	  attrs_found = true;
-	  break;
-	}
-
-      if (id & GOMP_TARGET_ARG_SUBSEQUENT_PARAM)
-	input++;
-      input++;
-    }
-
-  if (!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;
-      HSA_DEBUG ("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);
-
-  HSA_DEBUG ("GOMP_OFFLOAD_run called with %u dimensions:\n", kla->ndim);
-  unsigned i;
-  for (i = 0; i < kla->ndim; i++)
-    {
-      HSA_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;
-}
-
-/* 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;
-  struct agent_info *agent = kernel->agent;
-  create_and_finalize_hsa_program (agent);
-
-  if (agent->prog_finalized_error)
-    goto failure;
-
-  init_kernel (kernel);
-  if (kernel->initialization_failed)
-    goto failure;
-
-  return true;
-
-failure:
-  if (suppress_host_fallback)
-    GOMP_PLUGIN_fatal ("HSA host fallback has been suppressed");
-  HSA_DEBUG ("HSA target cannot be launched, doing a host fallback\n");
-  return false;
-}
-
-/* Atomically store pair of uint16_t values (HEADER and REST) to a PACKET.  */
-
-void
-packet_store_release (uint32_t* packet, uint16_t header, uint16_t rest)
-{
-  __atomic_store_n (packet, header | (rest << 16), __ATOMIC_RELEASE);
-}
-
-/* Run KERNEL on its agent, pass VARS to it as arguments and take
-   launchattributes from KLA.  */
-
-void
-run_kernel (struct kernel_info *kernel, void *vars,
-	    struct GOMP_kernel_launch_attributes *kla)
-{
-  struct agent_info *agent = kernel->agent;
-  if (pthread_rwlock_rdlock (&agent->modules_rwlock))
-    GOMP_PLUGIN_fatal ("Unable to read-lock an HSA agent rwlock");
-
-  if (!agent->initialized)
-    GOMP_PLUGIN_fatal ("Agent must be initialized");
-
-  if (!kernel->initialized)
-    GOMP_PLUGIN_fatal ("Called kernel must be initialized");
-
-  struct GOMP_hsa_kernel_dispatch *shadow
-    = create_kernel_dispatch (kernel, kernel->max_omp_data_size);
-
-  if (debug)
-    {
-      fprintf (stderr, "\nKernel has following dependencies:\n");
-      print_kernel_dispatch (shadow, 2);
-    }
-
-  uint64_t index
-    = hsa_fns.hsa_queue_add_write_index_release_fn (agent->command_q, 1);
-  HSA_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 (agent->command_q)
-	 >= agent->command_q->size)
-    ;
-
-  hsa_kernel_dispatch_packet_t *packet;
-  packet = ((hsa_kernel_dispatch_packet_t *) agent->command_q->base_address)
-	   + index % agent->command_q->size;
-
-  memset (((uint8_t *) packet) + 4, 0, sizeof (*packet) - 4);
-  packet->grid_size_x = kla->gdims[0];
-  packet->workgroup_size_x = get_group_size (kla->ndim, kla->gdims[0],
-					     kla->wdims[0]);
-
-  if (kla->ndim >= 2)
-    {
-      packet->grid_size_y = kla->gdims[1];
-      packet->workgroup_size_y = get_group_size (kla->ndim, kla->gdims[1],
-						 kla->wdims[1]);
-    }
-  else
-    {
-      packet->grid_size_y = 1;
-      packet->workgroup_size_y = 1;
-    }
-
-  if (kla->ndim == 3)
-    {
-      packet->grid_size_z = kla->gdims[2];
-      packet->workgroup_size_z = get_group_size (kla->ndim, kla->gdims[2],
-					     kla->wdims[2]);
-    }
-  else
-    {
-      packet->grid_size_z = 1;
-      packet->workgroup_size_z = 1;
-    }
-
-  packet->private_segment_size = kernel->private_segment_size;
-  packet->group_segment_size = kernel->group_segment_size;
-  packet->kernel_object = kernel->object;
-  packet->kernarg_address = shadow->kernarg_address;
-  hsa_signal_t s;
-  s.handle = shadow->signal;
-  packet->completion_signal = s;
-  hsa_fns.hsa_signal_store_relaxed_fn (s, 1);
-  memcpy (shadow->kernarg_address, &vars, sizeof (vars));
-
-  /* PR hsa/70337.  */
-  size_t vars_size = sizeof (vars);
-  if (kernel->kernarg_segment_size > vars_size)
-    {
-      if (kernel->kernarg_segment_size != vars_size
-	  + sizeof (struct hsa_kernel_runtime *))
-	GOMP_PLUGIN_fatal ("Kernel segment size has an unexpected value");
-      memcpy (packet->kernarg_address + vars_size, &shadow,
-	      sizeof (struct hsa_kernel_runtime *));
-    }
-
-  HSA_DEBUG ("Copying kernel runtime pointer to kernarg_address\n");
-
-  uint16_t header;
-  header = HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE;
-  header |= HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE;
-  header |= HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE;
-
-  HSA_DEBUG ("Going to dispatch kernel %s\n", kernel->name);
-
-  packet_store_release ((uint32_t *) packet, header,
-			(uint16_t) kla->ndim << HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS);
-
-  hsa_fns.hsa_signal_store_release_fn (agent->command_q->doorbell_signal,
-				       index);
-
-  /* TODO: GPU agents in Carrizo APUs cannot properly update L2 cache for
-     signal wait and signal load operations on their own and we need to
-     periodically call the hsa_signal_load_acquire on completion signals of
-     children kernels in the CPU to make that happen.  As soon the
-     limitation will be resolved, this workaround can be removed.  */
-
-  HSA_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)
-    for (unsigned i = 0; i < shadow->kernel_dispatch_count; i++)
-      {
-	hsa_signal_t child_s;
-	child_s.handle = shadow->children_dispatches[i]->signal;
-
-	HSA_DEBUG ("Waiting for children completion signal: %" PRIu64 "\n",
-		   (print_uint64_t) shadow->children_dispatches[i]->signal);
-	hsa_fns.hsa_signal_load_acquire_fn (child_s);
-      }
-
-  release_kernel_dispatch (shadow);
-
-  if (pthread_rwlock_unlock (&agent->modules_rwlock))
-    GOMP_PLUGIN_fatal ("Unable to unlock an HSA agent rwlock");
-}
-
-/* Part of the libgomp plugin interface.  Run a kernel on device N (the number
-   is actually ignored, we assume the FN_PTR has been mapped using the correct
-   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.  */
-
-void
-GOMP_OFFLOAD_run (int n __attribute__((unused)),
-		  void *fn_ptr, void *vars, void **args)
-{
-  struct kernel_info *kernel = (struct kernel_info *) fn_ptr;
-  struct GOMP_kernel_launch_attributes def;
-  struct GOMP_kernel_launch_attributes *kla;
-  if (!parse_target_attributes (args, &def, &kla))
-    {
-      HSA_DEBUG ("Will not run HSA kernel because the grid size is zero\n");
-      return;
-    }
-  run_kernel (kernel, vars, kla);
-}
-
-/* Information to be passed to a thread running a kernel asycnronously.  */
-
-struct async_run_info
-{
-  int device;
-  void *tgt_fn;
-  void *tgt_vars;
-  void **args;
-  void *async_data;
-};
-
-/* Thread routine to run a kernel asynchronously.  */
-
-static void *
-run_kernel_asynchronously (void *thread_arg)
-{
-  struct async_run_info *info = (struct async_run_info *) thread_arg;
-  int device = info->device;
-  void *tgt_fn = info->tgt_fn;
-  void *tgt_vars = info->tgt_vars;
-  void **args = info->args;
-  void *async_data = info->async_data;
-
-  free (info);
-  GOMP_OFFLOAD_run (device, tgt_fn, tgt_vars, args);
-  GOMP_PLUGIN_target_task_completion (async_data);
-  return NULL;
-}
-
-/* Part of the libgomp plugin interface.  Run a kernel like GOMP_OFFLOAD_run
-   does, but asynchronously and call GOMP_PLUGIN_target_task_completion when it
-   has finished.  */
-
-void
-GOMP_OFFLOAD_async_run (int device, void *tgt_fn, void *tgt_vars,
-			void **args, void *async_data)
-{
-  pthread_t pt;
-  struct async_run_info *info;
-  HSA_DEBUG ("GOMP_OFFLOAD_async_run invoked\n");
-  info = GOMP_PLUGIN_malloc (sizeof (struct async_run_info));
-
-  info->device = device;
-  info->tgt_fn = tgt_fn;
-  info->tgt_vars = tgt_vars;
-  info->args = args;
-  info->async_data = async_data;
-
-  int err = pthread_create (&pt, NULL, &run_kernel_asynchronously, info);
-  if (err != 0)
-    GOMP_PLUGIN_fatal ("HSA asynchronous thread creation failed: %s",
-		       strerror (err));
-  err = pthread_detach (pt);
-  if (err != 0)
-    GOMP_PLUGIN_fatal ("Failed to detach a thread to run HSA kernel "
-		       "asynchronously: %s", strerror (err));
-}
-
-/* Deinitialize all information associated with MODULE and kernels within
-   it.  Return TRUE on success.  */
-
-static bool
-destroy_module (struct module_info *module)
-{
-  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 an HSA kernel initialization "
-			   "mutex");
-	return false;
-      }
-  return true;
-}
-
-/* Part of the libgomp plugin interface.  Unload BRIG module described by
-   struct brig_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_HSA)
-    {
-      GOMP_PLUGIN_error ("Offload data incompatible with HSA plugin"
-			 " (expected %u, received %u)",
-			 GOMP_VERSION_HSA, GOMP_VERSION_DEV (version));
-      return false;
-    }
-
-  struct agent_info *agent;
-  agent = get_agent_info (n);
-  if (!agent)
-    return false;
-
-  if (pthread_rwlock_wrlock (&agent->modules_rwlock))
-    {
-      GOMP_PLUGIN_error ("Unable to write-lock an HSA agent rwlock");
-      return false;
-    }
-  struct module_info *module = agent->first_module;
-  while (module)
-    {
-      if (module->image_desc == target_data)
-	break;
-      module = module->next;
-    }
-  if (!module)
-    {
-      GOMP_PLUGIN_error ("Attempt to unload an image that has never been "
-			 "loaded before");
-      return false;
-    }
-
-  remove_module_from_agent (agent, module);
-  if (!destroy_module (module))
-    return false;
-  free (module);
-  if (!destroy_hsa_program (agent))
-    return false;
-  if (pthread_rwlock_unlock (&agent->modules_rwlock))
-    {
-      GOMP_PLUGIN_error ("Unable to unlock an HSA agent rwlock");
-      return false;
-    }
-  return true;
-}
-
-/* Part of the libgomp plugin interface.  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;
-
-  struct module_info *next_module = agent->first_module;
-  while (next_module)
-    {
-      struct module_info *module = next_module;
-      next_module = module->next;
-      if (!destroy_module (module))
-	return false;
-      free (module);
-    }
-  agent->first_module = NULL;
-  if (!destroy_hsa_program (agent))
-    return false;
-
-  release_agent_shared_libraries (agent);
-
-  hsa_status_t status = hsa_fns.hsa_queue_destroy_fn (agent->command_q);
-  if (status != HSA_STATUS_SUCCESS)
-    return hsa_error ("Error destroying command queue", status);
-  status = hsa_fns.hsa_queue_destroy_fn (agent->kernel_dispatch_command_q);
-  if (status != HSA_STATUS_SUCCESS)
-    return hsa_error ("Error destroying kernel dispatch command queue", status);
-  if (pthread_mutex_destroy (&agent->prog_mutex))
-    {
-      GOMP_PLUGIN_error ("Failed to destroy an HSA agent program mutex");
-      return false;
-    }
-  if (pthread_rwlock_destroy (&agent->modules_rwlock))
-    {
-      GOMP_PLUGIN_error ("Failed to destroy an HSA agent rwlock");
-      return false;
-    }
-  agent->initialized = false;
-  return true;
-}
-
-/* Part of the libgomp plugin interface.  Not implemented as it is not required
-   for HSA.  */
-
-void *
-GOMP_OFFLOAD_alloc (int ord, size_t size)
-{
-  GOMP_PLUGIN_error ("HSA GOMP_OFFLOAD_alloc is not implemented because "
-		     "it should never be called");
-  return NULL;
-}
-
-/* Part of the libgomp plugin interface.  Not implemented as it is not required
-   for HSA.  */
-
-bool
-GOMP_OFFLOAD_free (int ord, void *ptr)
-{
-  GOMP_PLUGIN_error ("HSA GOMP_OFFLOAD_free is not implemented because "
-		     "it should never be called");
-  return false;
-}
-
-/* Part of the libgomp plugin interface.  Not implemented as it is not required
-   for HSA.  */
-
-bool
-GOMP_OFFLOAD_dev2host (int ord, void *dst, const void *src, size_t n)
-{
-  GOMP_PLUGIN_error ("HSA GOMP_OFFLOAD_dev2host is not implemented because "
-		     "it should never be called");
-  return false;
-}
-
-/* Part of the libgomp plugin interface.  Not implemented as it is not required
-   for HSA.  */
-
-bool
-GOMP_OFFLOAD_host2dev (int ord, void *dst, const void *src, size_t n)
-{
-  GOMP_PLUGIN_error ("HSA GOMP_OFFLOAD_host2dev is not implemented because "
-		     "it should never be called");
-  return false;
-}
-
-/* Part of the libgomp plugin interface.  Not implemented as it is not required
-   for HSA.  */
-
-bool
-GOMP_OFFLOAD_dev2dev (int ord, void *dst, const void *src, size_t n)
-{
-  GOMP_PLUGIN_error ("HSA GOMP_OFFLOAD_dev2dev is not implemented because "
-		     "it should never be called");
-  return false;
-}