+2016-11-23 Martin Liska <mliska@suse.cz>
+ Martin Jambor <mjambor@suse.cz>
+
+ * doc/install.texi: Remove entry about --with-hsa-kmt-lib.
+
2016-11-23 Aldy Hernandez <aldyh@redhat.com>
PR target/78213
shorthand for
@option{--with-hsa-runtime-lib=@/@var{hsainstalldir}/lib} and
@option{--with-hsa-runtime-include=@/@var{hsainstalldir}/include}.
-
-@item --with-hsa-kmt-lib=@var{pathname}
-
-If you configure GCC with HSA offloading but do not have the HSA
-KMT library installed in a standard location then you can
-explicitly specify the directory where it resides.
@end table
@subheading Cross-Compiler-Specific Options
+2016-11-23 Martin Liska <mliska@suse.cz>
+ Martin Jambor <mjambor@suse.cz>
+
+ * plugin/hsa.h: New file.
+ * plugin/hsa_ext_finalize.h: New file.
+ * plugin/configfrag.ac: Remove hsa-kmt-lib test. Added checks for
+ header file unistd.h, and functions secure_getenv, __secure_getenv,
+ getuid, geteuid, getgid and getegid.
+ * plugin/Makefrag.am (libgomp_plugin_hsa_la_CPPFLAGS): Added
+ -D_GNU_SOURCE.
+ * plugin/plugin-hsa.c: Include config.h, inttypes.h and stdbool.h.
+ Handle various cases of secure_getenv presence, add an implementation
+ when we can test effective UID and GID.
+ (struct hsa_runtime_fn_info): New structure.
+ (hsa_runtime_fn_info hsa_fns): New variable.
+ (hsa_runtime_lib): Likewise.
+ (support_cpu_devices): Likewise.
+ (init_enviroment_variables): Load newly introduced ENV
+ variables.
+ (hsa_warn): Call hsa run-time functions via hsa_fns structure.
+ (hsa_fatal): Likewise.
+ (DLSYM_FN): New macro.
+ (init_hsa_runtime_functions): New function.
+ (suitable_hsa_agent_p): Call hsa run-time functions via hsa_fns
+ structure. Depending on environment, also allow CPU devices.
+ (init_hsa_context): Call hsa run-time functions via hsa_fns structure.
+ (get_kernarg_memory_region): Likewise.
+ (GOMP_OFFLOAD_init_device): Likewise.
+ (destroy_hsa_program): Likewise.
+ (init_basic_kernel_info): New function.
+ (GOMP_OFFLOAD_load_image): Use it.
+ (create_and_finalize_hsa_program): Call hsa run-time functions via
+ hsa_fns structure.
+ (create_single_kernel_dispatch): Likewise.
+ (release_kernel_dispatch): Likewise.
+ (init_single_kernel): Likewise.
+ (parse_target_attributes): Allow up multiple HSA grid dimensions.
+ (get_group_size): New function.
+ (run_kernel): Likewise.
+ (GOMP_OFFLOAD_run): Outline most functionality to run_kernel.
+ (GOMP_OFFLOAD_fini_device): Call hsa run-time functions via hsa_fns
+ structure.
+ * testsuite/lib/libgomp.exp: Remove hsa_kmt_lib support.
+ * testsuite/libgomp-test-support.exp.in: Likewise.
+ * Makefile.in: Regenerated.
+ * aclocal.m4: Likewise.
+ * config.h.in: Likewise.
+ * configure: Likewise.
+ * testsuite/Makefile.in: Likewise.
+
2016-11-15 Martin Jambor <mjambor@suse.cz>
Alexander Monakov <amonakov@ispras.ru>
-# Makefile.in generated by automake 1.11.6 from Makefile.am.
+# Makefile.in generated by automake 1.11.1 from Makefile.am.
# @configure_input@
# Copyright (C) 1994, 1995, 1996, 1997, 1998, 1999, 2000, 2001, 2002,
-# 2003, 2004, 2005, 2006, 2007, 2008, 2009, 2010, 2011 Free Software
-# Foundation, Inc.
+# 2003, 2004, 2005, 2006, 2007, 2008, 2009 Free Software Foundation,
+# Inc.
# This Makefile.in is free software; the Free Software Foundation
# gives unlimited permission to copy and/or distribute it,
# with or without modifications, as long as this notice is preserved.
VPATH = @srcdir@
-am__make_dryrun = \
- { \
- am__dry=no; \
- case $$MAKEFLAGS in \
- *\\[\ \ ]*) \
- echo 'am--echo: ; @echo "AM" OK' | $(MAKE) -f - 2>/dev/null \
- | grep '^AM OK$$' >/dev/null || am__dry=yes;; \
- *) \
- for am__flg in $$MAKEFLAGS; do \
- case $$am__flg in \
- *=*|--*) ;; \
- *n*) am__dry=yes; break;; \
- esac; \
- done;; \
- esac; \
- test $$am__dry = yes; \
- }
pkgdatadir = $(datadir)/@PACKAGE@
pkgincludedir = $(includedir)/@PACKAGE@
pkglibdir = $(libdir)/@PACKAGE@
am__base_list = \
sed '$$!N;$$!N;$$!N;$$!N;$$!N;$$!N;$$!N;s/\n/ /g' | \
sed '$$!N;$$!N;$$!N;$$!N;s/\n/ /g'
-am__uninstall_files_from_dir = { \
- test -z "$$files" \
- || { test ! -d "$$dir" && test ! -f "$$dir" && test ! -r "$$dir"; } \
- || { echo " ( cd '$$dir' && rm -f" $$files ")"; \
- $(am__cd) "$$dir" && rm -f $$files; }; \
- }
am__installdirs = "$(DESTDIR)$(toolexeclibdir)" "$(DESTDIR)$(infodir)" \
"$(DESTDIR)$(fincludedir)" "$(DESTDIR)$(libsubincludedir)" \
"$(DESTDIR)$(toolexeclibdir)"
install-pdf-recursive install-ps-recursive install-recursive \
installcheck-recursive installdirs-recursive pdf-recursive \
ps-recursive uninstall-recursive
-am__can_run_installinfo = \
- case $$AM_UPDATE_INFO_DIR in \
- n|no|NO) false;; \
- *) (install-info --version) >/dev/null 2>&1;; \
- esac
HEADERS = $(nodist_finclude_HEADERS) $(nodist_libsubinclude_HEADERS) \
$(nodist_noinst_HEADERS) $(nodist_toolexeclib_HEADERS)
RECURSIVE_CLEAN_TARGETS = mostlyclean-recursive clean-recursive \
FCFLAGS = @FCFLAGS@
FGREP = @FGREP@
GREP = @GREP@
-HSA_KMT_LIB = @HSA_KMT_LIB@
HSA_RUNTIME_INCLUDE = @HSA_RUNTIME_INCLUDE@
HSA_RUNTIME_LIB = @HSA_RUNTIME_LIB@
INSTALL = @INSTALL@
# Heterogenous Systems Architecture plugin
@PLUGIN_HSA_TRUE@libgomp_plugin_hsa_version_info = -version-info $(libtool_VERSION)
@PLUGIN_HSA_TRUE@libgomp_plugin_hsa_la_SOURCES = plugin/plugin-hsa.c
-@PLUGIN_HSA_TRUE@libgomp_plugin_hsa_la_CPPFLAGS = $(AM_CPPFLAGS) $(PLUGIN_HSA_CPPFLAGS)
+@PLUGIN_HSA_TRUE@libgomp_plugin_hsa_la_CPPFLAGS = $(AM_CPPFLAGS) $(PLUGIN_HSA_CPPFLAGS) \
+@PLUGIN_HSA_TRUE@ -D_GNU_SOURCE
+
@PLUGIN_HSA_TRUE@libgomp_plugin_hsa_la_LDFLAGS = \
@PLUGIN_HSA_TRUE@ $(libgomp_plugin_hsa_version_info) \
@PLUGIN_HSA_TRUE@ $(lt_host_flags) $(PLUGIN_HSA_LDFLAGS)
.SUFFIXES:
.SUFFIXES: .c .dvi .f90 .lo .o .obj .ps
-am--refresh: Makefile
+am--refresh:
@:
$(srcdir)/Makefile.in: @MAINTAINER_MODE_TRUE@ $(srcdir)/Makefile.am $(top_srcdir)/plugin/Makefrag.am $(am__configure_deps)
@for dep in $?; do \
echo ' cd $(top_builddir) && $(SHELL) ./config.status $@ $(am__depfiles_maybe)'; \
cd $(top_builddir) && $(SHELL) ./config.status $@ $(am__depfiles_maybe);; \
esac;
-$(top_srcdir)/plugin/Makefrag.am:
$(top_builddir)/config.status: $(top_srcdir)/configure $(CONFIG_STATUS_DEPENDENCIES)
$(SHELL) ./config.status --recheck
$(am__aclocal_m4_deps):
config.h: stamp-h1
- @if test ! -f $@; then rm -f stamp-h1; else :; fi
- @if test ! -f $@; then $(MAKE) $(AM_MAKEFLAGS) stamp-h1; else :; fi
+ @if test ! -f $@; then \
+ rm -f stamp-h1; \
+ $(MAKE) $(AM_MAKEFLAGS) stamp-h1; \
+ else :; fi
stamp-h1: $(srcdir)/config.h.in $(top_builddir)/config.status
@rm -f stamp-h1
cd $(top_builddir) && $(SHELL) ./config.status $@
install-toolexeclibLTLIBRARIES: $(toolexeclib_LTLIBRARIES)
@$(NORMAL_INSTALL)
+ test -z "$(toolexeclibdir)" || $(MKDIR_P) "$(DESTDIR)$(toolexeclibdir)"
@list='$(toolexeclib_LTLIBRARIES)'; test -n "$(toolexeclibdir)" || list=; \
list2=; for p in $$list; do \
if test -f $$p; then \
else :; fi; \
done; \
test -z "$$list2" || { \
- echo " $(MKDIR_P) '$(DESTDIR)$(toolexeclibdir)'"; \
- $(MKDIR_P) "$(DESTDIR)$(toolexeclibdir)" || exit 1; \
echo " $(LIBTOOL) $(AM_LIBTOOLFLAGS) $(LIBTOOLFLAGS) --mode=install $(INSTALL) $(INSTALL_STRIP_FLAG) $$list2 '$(DESTDIR)$(toolexeclibdir)'"; \
$(LIBTOOL) $(AM_LIBTOOLFLAGS) $(LIBTOOLFLAGS) --mode=install $(INSTALL) $(INSTALL_STRIP_FLAG) $$list2 "$(DESTDIR)$(toolexeclibdir)"; \
}
echo "rm -f \"$${dir}/so_locations\""; \
rm -f "$${dir}/so_locations"; \
done
-libgomp-plugin-hsa.la: $(libgomp_plugin_hsa_la_OBJECTS) $(libgomp_plugin_hsa_la_DEPENDENCIES) $(EXTRA_libgomp_plugin_hsa_la_DEPENDENCIES)
+libgomp-plugin-hsa.la: $(libgomp_plugin_hsa_la_OBJECTS) $(libgomp_plugin_hsa_la_DEPENDENCIES)
$(libgomp_plugin_hsa_la_LINK) $(am_libgomp_plugin_hsa_la_rpath) $(libgomp_plugin_hsa_la_OBJECTS) $(libgomp_plugin_hsa_la_LIBADD) $(LIBS)
-libgomp-plugin-nvptx.la: $(libgomp_plugin_nvptx_la_OBJECTS) $(libgomp_plugin_nvptx_la_DEPENDENCIES) $(EXTRA_libgomp_plugin_nvptx_la_DEPENDENCIES)
+libgomp-plugin-nvptx.la: $(libgomp_plugin_nvptx_la_OBJECTS) $(libgomp_plugin_nvptx_la_DEPENDENCIES)
$(libgomp_plugin_nvptx_la_LINK) $(am_libgomp_plugin_nvptx_la_rpath) $(libgomp_plugin_nvptx_la_OBJECTS) $(libgomp_plugin_nvptx_la_LIBADD) $(LIBS)
-libgomp.la: $(libgomp_la_OBJECTS) $(libgomp_la_DEPENDENCIES) $(EXTRA_libgomp_la_DEPENDENCIES)
+libgomp.la: $(libgomp_la_OBJECTS) $(libgomp_la_DEPENDENCIES)
$(libgomp_la_LINK) -rpath $(toolexeclibdir) $(libgomp_la_OBJECTS) $(libgomp_la_LIBADD) $(LIBS)
mostlyclean-compile:
uninstall-info-am:
@$(PRE_UNINSTALL)
- @if test -d '$(DESTDIR)$(infodir)' && $(am__can_run_installinfo); then \
+ @if test -d '$(DESTDIR)$(infodir)' && \
+ (install-info --version && \
+ install-info --version 2>&1 | sed 1q | grep -i -v debian) >/dev/null 2>&1; then \
list='$(INFO_DEPS)'; \
for file in $$list; do \
relfile=`echo "$$file" | sed 's|^.*/||'`; \
done
install-nodist_fincludeHEADERS: $(nodist_finclude_HEADERS)
@$(NORMAL_INSTALL)
+ test -z "$(fincludedir)" || $(MKDIR_P) "$(DESTDIR)$(fincludedir)"
@list='$(nodist_finclude_HEADERS)'; test -n "$(fincludedir)" || list=; \
- if test -n "$$list"; then \
- echo " $(MKDIR_P) '$(DESTDIR)$(fincludedir)'"; \
- $(MKDIR_P) "$(DESTDIR)$(fincludedir)" || exit 1; \
- fi; \
for p in $$list; do \
if test -f "$$p"; then d=; else d="$(srcdir)/"; fi; \
echo "$$d$$p"; \
@$(NORMAL_UNINSTALL)
@list='$(nodist_finclude_HEADERS)'; test -n "$(fincludedir)" || list=; \
files=`for p in $$list; do echo $$p; done | sed -e 's|^.*/||'`; \
- dir='$(DESTDIR)$(fincludedir)'; $(am__uninstall_files_from_dir)
+ test -n "$$files" || exit 0; \
+ echo " ( cd '$(DESTDIR)$(fincludedir)' && rm -f" $$files ")"; \
+ cd "$(DESTDIR)$(fincludedir)" && rm -f $$files
install-nodist_libsubincludeHEADERS: $(nodist_libsubinclude_HEADERS)
@$(NORMAL_INSTALL)
+ test -z "$(libsubincludedir)" || $(MKDIR_P) "$(DESTDIR)$(libsubincludedir)"
@list='$(nodist_libsubinclude_HEADERS)'; test -n "$(libsubincludedir)" || list=; \
- if test -n "$$list"; then \
- echo " $(MKDIR_P) '$(DESTDIR)$(libsubincludedir)'"; \
- $(MKDIR_P) "$(DESTDIR)$(libsubincludedir)" || exit 1; \
- fi; \
for p in $$list; do \
if test -f "$$p"; then d=; else d="$(srcdir)/"; fi; \
echo "$$d$$p"; \
@$(NORMAL_UNINSTALL)
@list='$(nodist_libsubinclude_HEADERS)'; test -n "$(libsubincludedir)" || list=; \
files=`for p in $$list; do echo $$p; done | sed -e 's|^.*/||'`; \
- dir='$(DESTDIR)$(libsubincludedir)'; $(am__uninstall_files_from_dir)
+ test -n "$$files" || exit 0; \
+ echo " ( cd '$(DESTDIR)$(libsubincludedir)' && rm -f" $$files ")"; \
+ cd "$(DESTDIR)$(libsubincludedir)" && rm -f $$files
install-nodist_toolexeclibHEADERS: $(nodist_toolexeclib_HEADERS)
@$(NORMAL_INSTALL)
+ test -z "$(toolexeclibdir)" || $(MKDIR_P) "$(DESTDIR)$(toolexeclibdir)"
@list='$(nodist_toolexeclib_HEADERS)'; test -n "$(toolexeclibdir)" || list=; \
- if test -n "$$list"; then \
- echo " $(MKDIR_P) '$(DESTDIR)$(toolexeclibdir)'"; \
- $(MKDIR_P) "$(DESTDIR)$(toolexeclibdir)" || exit 1; \
- fi; \
for p in $$list; do \
if test -f "$$p"; then d=; else d="$(srcdir)/"; fi; \
echo "$$d$$p"; \
@$(NORMAL_UNINSTALL)
@list='$(nodist_toolexeclib_HEADERS)'; test -n "$(toolexeclibdir)" || list=; \
files=`for p in $$list; do echo $$p; done | sed -e 's|^.*/||'`; \
- dir='$(DESTDIR)$(toolexeclibdir)'; $(am__uninstall_files_from_dir)
+ test -n "$$files" || exit 0; \
+ echo " ( cd '$(DESTDIR)$(toolexeclibdir)' && rm -f" $$files ")"; \
+ cd "$(DESTDIR)$(toolexeclibdir)" && rm -f $$files
# This directory's subdirectories are mostly independent; you can cd
# into them and run `make' without going through this Makefile.
installcheck: installcheck-recursive
install-strip:
- if test -z '$(STRIP)'; then \
- $(MAKE) $(AM_MAKEFLAGS) INSTALL_PROGRAM="$(INSTALL_STRIP_PROGRAM)" \
- install_sh_PROGRAM="$(INSTALL_STRIP_PROGRAM)" INSTALL_STRIP_FLAG=-s \
- install; \
- else \
- $(MAKE) $(AM_MAKEFLAGS) INSTALL_PROGRAM="$(INSTALL_STRIP_PROGRAM)" \
- install_sh_PROGRAM="$(INSTALL_STRIP_PROGRAM)" INSTALL_STRIP_FLAG=-s \
- "INSTALL_PROGRAM_ENV=STRIPPROG='$(STRIP)'" install; \
- fi
+ $(MAKE) $(AM_MAKEFLAGS) INSTALL_PROGRAM="$(INSTALL_STRIP_PROGRAM)" \
+ install_sh_PROGRAM="$(INSTALL_STRIP_PROGRAM)" INSTALL_STRIP_FLAG=-s \
+ `test -z '$(STRIP)' || \
+ echo "INSTALL_PROGRAM_ENV=STRIPPROG='$(STRIP)'"` install
mostlyclean-generic:
clean-generic:
install-dvi-am: $(DVIS)
@$(NORMAL_INSTALL)
+ test -z "$(dvidir)" || $(MKDIR_P) "$(DESTDIR)$(dvidir)"
@list='$(DVIS)'; test -n "$(dvidir)" || list=; \
- if test -n "$$list"; then \
- echo " $(MKDIR_P) '$(DESTDIR)$(dvidir)'"; \
- $(MKDIR_P) "$(DESTDIR)$(dvidir)" || exit 1; \
- fi; \
for p in $$list; do \
if test -f "$$p"; then d=; else d="$(srcdir)/"; fi; \
echo "$$d$$p"; \
install-html-am: $(HTMLS)
@$(NORMAL_INSTALL)
+ test -z "$(htmldir)" || $(MKDIR_P) "$(DESTDIR)$(htmldir)"
@list='$(HTMLS)'; list2=; test -n "$(htmldir)" || list=; \
- if test -n "$$list"; then \
- echo " $(MKDIR_P) '$(DESTDIR)$(htmldir)'"; \
- $(MKDIR_P) "$(DESTDIR)$(htmldir)" || exit 1; \
- fi; \
for p in $$list; do \
if test -f "$$p" || test -d "$$p"; then d=; else d="$(srcdir)/"; fi; \
$(am__strip_dir) \
- d2=$$d$$p; \
- if test -d "$$d2"; then \
+ if test -d "$$d$$p"; then \
echo " $(MKDIR_P) '$(DESTDIR)$(htmldir)/$$f'"; \
$(MKDIR_P) "$(DESTDIR)$(htmldir)/$$f" || exit 1; \
- echo " $(INSTALL_DATA) '$$d2'/* '$(DESTDIR)$(htmldir)/$$f'"; \
- $(INSTALL_DATA) "$$d2"/* "$(DESTDIR)$(htmldir)/$$f" || exit $$?; \
+ echo " $(INSTALL_DATA) '$$d$$p'/* '$(DESTDIR)$(htmldir)/$$f'"; \
+ $(INSTALL_DATA) "$$d$$p"/* "$(DESTDIR)$(htmldir)/$$f" || exit $$?; \
else \
- list2="$$list2 $$d2"; \
+ list2="$$list2 $$d$$p"; \
fi; \
done; \
test -z "$$list2" || { echo "$$list2" | $(am__base_list) | \
install-info-am: $(INFO_DEPS)
@$(NORMAL_INSTALL)
+ test -z "$(infodir)" || $(MKDIR_P) "$(DESTDIR)$(infodir)"
@srcdirstrip=`echo "$(srcdir)" | sed 's|.|.|g'`; \
list='$(INFO_DEPS)'; test -n "$(infodir)" || list=; \
- if test -n "$$list"; then \
- echo " $(MKDIR_P) '$(DESTDIR)$(infodir)'"; \
- $(MKDIR_P) "$(DESTDIR)$(infodir)" || exit 1; \
- fi; \
for file in $$list; do \
case $$file in \
$(srcdir)/*) file=`echo "$$file" | sed "s|^$$srcdirstrip/||"`;; \
echo " $(INSTALL_DATA) $$files '$(DESTDIR)$(infodir)'"; \
$(INSTALL_DATA) $$files "$(DESTDIR)$(infodir)" || exit $$?; done
@$(POST_INSTALL)
- @if $(am__can_run_installinfo); then \
+ @if (install-info --version && \
+ install-info --version 2>&1 | sed 1q | grep -i -v debian) >/dev/null 2>&1; then \
list='$(INFO_DEPS)'; test -n "$(infodir)" || list=; \
for file in $$list; do \
relfile=`echo "$$file" | sed 's|^.*/||'`; \
install-pdf-am: $(PDFS)
@$(NORMAL_INSTALL)
+ test -z "$(pdfdir)" || $(MKDIR_P) "$(DESTDIR)$(pdfdir)"
@list='$(PDFS)'; test -n "$(pdfdir)" || list=; \
- if test -n "$$list"; then \
- echo " $(MKDIR_P) '$(DESTDIR)$(pdfdir)'"; \
- $(MKDIR_P) "$(DESTDIR)$(pdfdir)" || exit 1; \
- fi; \
for p in $$list; do \
if test -f "$$p"; then d=; else d="$(srcdir)/"; fi; \
echo "$$d$$p"; \
install-ps-am: $(PSS)
@$(NORMAL_INSTALL)
+ test -z "$(psdir)" || $(MKDIR_P) "$(DESTDIR)$(psdir)"
@list='$(PSS)'; test -n "$(psdir)" || list=; \
- if test -n "$$list"; then \
- echo " $(MKDIR_P) '$(DESTDIR)$(psdir)'"; \
- $(MKDIR_P) "$(DESTDIR)$(psdir)" || exit 1; \
- fi; \
for p in $$list; do \
if test -f "$$p"; then d=; else d="$(srcdir)/"; fi; \
echo "$$d$$p"; \
-# generated automatically by aclocal 1.11.6 -*- Autoconf -*-
+# generated automatically by aclocal 1.11.1 -*- Autoconf -*-
# Copyright (C) 1996, 1997, 1998, 1999, 2000, 2001, 2002, 2003, 2004,
-# 2005, 2006, 2007, 2008, 2009, 2010, 2011 Free Software Foundation,
-# Inc.
+# 2005, 2006, 2007, 2008, 2009 Free Software Foundation, Inc.
# This file is free software; the Free Software Foundation
# gives unlimited permission to copy and/or distribute it,
# with or without modifications, as long as this notice is preserved.
If you have problems, you may need to regenerate the build system entirely.
To do so, use the procedure documented by the package, typically `autoreconf'.])])
-# Copyright (C) 2002, 2003, 2005, 2006, 2007, 2008, 2011 Free Software
-# Foundation, Inc.
+# Copyright (C) 2002, 2003, 2005, 2006, 2007, 2008 Free Software Foundation, Inc.
#
# This file is free software; the Free Software Foundation
# gives unlimited permission to copy and/or distribute it,
# with or without modifications, as long as this notice is preserved.
-# serial 1
-
# AM_AUTOMAKE_VERSION(VERSION)
# ----------------------------
# Automake X.Y traces this macro to ensure aclocal.m4 has been
[am__api_version='1.11'
dnl Some users find AM_AUTOMAKE_VERSION and mistake it for a way to
dnl require some minimum version. Point them to the right macro.
-m4_if([$1], [1.11.6], [],
+m4_if([$1], [1.11.1], [],
[AC_FATAL([Do not call $0, use AM_INIT_AUTOMAKE([$1]).])])dnl
])
# Call AM_AUTOMAKE_VERSION and AM_AUTOMAKE_VERSION so they can be traced.
# This function is AC_REQUIREd by AM_INIT_AUTOMAKE.
AC_DEFUN([AM_SET_CURRENT_AUTOMAKE_VERSION],
-[AM_AUTOMAKE_VERSION([1.11.6])dnl
+[AM_AUTOMAKE_VERSION([1.11.1])dnl
m4_ifndef([AC_AUTOCONF_VERSION],
[m4_copy([m4_PACKAGE_VERSION], [AC_AUTOCONF_VERSION])])dnl
_AM_AUTOCONF_VERSION(m4_defn([AC_AUTOCONF_VERSION]))])
# AM_AUX_DIR_EXPAND -*- Autoconf -*-
-# Copyright (C) 2001, 2003, 2005, 2011 Free Software Foundation, Inc.
+# Copyright (C) 2001, 2003, 2005 Free Software Foundation, Inc.
#
# This file is free software; the Free Software Foundation
# gives unlimited permission to copy and/or distribute it,
# with or without modifications, as long as this notice is preserved.
-# serial 1
-
# For projects using AC_CONFIG_AUX_DIR([foo]), Autoconf sets
# $ac_aux_dir to `$srcdir/foo'. In other projects, it is set to
# `$srcdir', `$srcdir/..', or `$srcdir/../..'.
Usually this means the macro was only invoked conditionally.]])
fi])])
-# Copyright (C) 1999, 2000, 2001, 2002, 2003, 2004, 2005, 2006, 2009,
-# 2010, 2011 Free Software Foundation, Inc.
+# Copyright (C) 1999, 2000, 2001, 2002, 2003, 2004, 2005, 2006, 2009
+# Free Software Foundation, Inc.
#
# This file is free software; the Free Software Foundation
# gives unlimited permission to copy and/or distribute it,
# with or without modifications, as long as this notice is preserved.
-# serial 12
+# serial 10
# There are a few dirty hacks below to avoid letting `AC_PROG_CC' be
# written in clear, in which case automake, when reading aclocal.m4,
# instance it was reported that on HP-UX the gcc test will end up
# making a dummy file named `D' -- because `-MD' means `put the output
# in D'.
- rm -rf conftest.dir
mkdir conftest.dir
# Copy depcomp to subdir because otherwise we won't find it if we're
# using a relative directory.
break
fi
;;
- msvc7 | msvc7msys | msvisualcpp | msvcmsys)
+ msvisualcpp | msvcmsys)
# This compiler won't grok `-c -o', but also, the minuso test has
# not run yet. These depmodes are late enough in the game, and
# so weak that their functioning should not be impacted.
if test "x$enable_dependency_tracking" != xno; then
am_depcomp="$ac_aux_dir/depcomp"
AMDEPBACKSLASH='\'
- am__nodep='_no'
fi
AM_CONDITIONAL([AMDEP], [test "x$enable_dependency_tracking" != xno])
AC_SUBST([AMDEPBACKSLASH])dnl
_AM_SUBST_NOTMAKE([AMDEPBACKSLASH])dnl
-AC_SUBST([am__nodep])dnl
-_AM_SUBST_NOTMAKE([am__nodep])dnl
])
# Generate code to set up dependency tracking. -*- Autoconf -*-
done
echo "timestamp for $_am_arg" >`AS_DIRNAME(["$_am_arg"])`/stamp-h[]$_am_stamp_count])
-# Copyright (C) 2001, 2003, 2005, 2008, 2011 Free Software Foundation,
-# Inc.
+# Copyright (C) 2001, 2003, 2005, 2008 Free Software Foundation, Inc.
#
# This file is free software; the Free Software Foundation
# gives unlimited permission to copy and/or distribute it,
# with or without modifications, as long as this notice is preserved.
-# serial 1
-
# AM_PROG_INSTALL_SH
# ------------------
# Define $install_sh.
# Add --enable-maintainer-mode option to configure. -*- Autoconf -*-
# From Jim Meyering
-# Copyright (C) 1996, 1998, 2000, 2001, 2002, 2003, 2004, 2005, 2008,
-# 2011 Free Software Foundation, Inc.
+# Copyright (C) 1996, 1998, 2000, 2001, 2002, 2003, 2004, 2005, 2008
+# Free Software Foundation, Inc.
#
# This file is free software; the Free Software Foundation
# gives unlimited permission to copy and/or distribute it,
[disable], [m4_define([am_maintainer_other], [enable])],
[m4_define([am_maintainer_other], [enable])
m4_warn([syntax], [unexpected argument to AM@&t@_MAINTAINER_MODE: $1])])
-AC_MSG_CHECKING([whether to enable maintainer-specific portions of Makefiles])
+AC_MSG_CHECKING([whether to am_maintainer_other maintainer-specific portions of Makefiles])
dnl maintainer-mode's default is 'disable' unless 'enable' is passed
AC_ARG_ENABLE([maintainer-mode],
[ --][am_maintainer_other][-maintainer-mode am_maintainer_other make rules and dependencies not useful
fi
])
-# Copyright (C) 2003, 2004, 2005, 2006, 2011 Free Software Foundation,
-# Inc.
+# Copyright (C) 2003, 2004, 2005, 2006 Free Software Foundation, Inc.
#
# This file is free software; the Free Software Foundation
# gives unlimited permission to copy and/or distribute it,
# with or without modifications, as long as this notice is preserved.
-# serial 1
-
# AM_PROG_MKDIR_P
# ---------------
# Check for `mkdir -p'.
# Helper functions for option handling. -*- Autoconf -*-
-# Copyright (C) 2001, 2002, 2003, 2005, 2008, 2010 Free Software
-# Foundation, Inc.
+# Copyright (C) 2001, 2002, 2003, 2005, 2008 Free Software Foundation, Inc.
#
# This file is free software; the Free Software Foundation
# gives unlimited permission to copy and/or distribute it,
# with or without modifications, as long as this notice is preserved.
-# serial 5
+# serial 4
# _AM_MANGLE_OPTION(NAME)
# -----------------------
[[_AM_OPTION_]m4_bpatsubst($1, [[^a-zA-Z0-9_]], [_])])
# _AM_SET_OPTION(NAME)
-# --------------------
+# ------------------------------
# Set option NAME. Presently that only means defining a flag for this option.
AC_DEFUN([_AM_SET_OPTION],
[m4_define(_AM_MANGLE_OPTION([$1]), 1)])
# _AM_SET_OPTIONS(OPTIONS)
-# ------------------------
+# ----------------------------------
# OPTIONS is a space-separated list of Automake options.
AC_DEFUN([_AM_SET_OPTIONS],
[m4_foreach_w([_AM_Option], [$1], [_AM_SET_OPTION(_AM_Option)])])
fi
AC_MSG_RESULT(yes)])
-# Copyright (C) 2001, 2003, 2005, 2011 Free Software Foundation, Inc.
+# Copyright (C) 2001, 2003, 2005 Free Software Foundation, Inc.
#
# This file is free software; the Free Software Foundation
# gives unlimited permission to copy and/or distribute it,
# with or without modifications, as long as this notice is preserved.
-# serial 1
-
# AM_PROG_INSTALL_STRIP
# ---------------------
# One issue with vendor `install' (even GNU) is that you can't
INSTALL_STRIP_PROGRAM="\$(install_sh) -c -s"
AC_SUBST([INSTALL_STRIP_PROGRAM])])
-# Copyright (C) 2006, 2008, 2010 Free Software Foundation, Inc.
+# Copyright (C) 2006, 2008 Free Software Foundation, Inc.
#
# This file is free software; the Free Software Foundation
# gives unlimited permission to copy and/or distribute it,
# with or without modifications, as long as this notice is preserved.
-# serial 3
+# serial 2
# _AM_SUBST_NOTMAKE(VARIABLE)
# ---------------------------
AC_DEFUN([_AM_SUBST_NOTMAKE])
# AM_SUBST_NOTMAKE(VARIABLE)
-# --------------------------
+# ---------------------------
# Public sister of _AM_SUBST_NOTMAKE.
AC_DEFUN([AM_SUBST_NOTMAKE], [_AM_SUBST_NOTMAKE($@)])
# Check how to create a tarball. -*- Autoconf -*-
-# Copyright (C) 2004, 2005, 2012 Free Software Foundation, Inc.
+# Copyright (C) 2004, 2005 Free Software Foundation, Inc.
#
# This file is free software; the Free Software Foundation
# gives unlimited permission to copy and/or distribute it,
# a tarball read from stdin.
# $(am__untar) < result.tar
AC_DEFUN([_AM_PROG_TAR],
-[# Always define AMTAR for backward compatibility. Yes, it's still used
-# in the wild :-( We should find a proper way to deprecate it ...
-AC_SUBST([AMTAR], ['$${TAR-tar}'])
+[# Always define AMTAR for backward compatibility.
+AM_MISSING_PROG([AMTAR], [tar])
m4_if([$1], [v7],
- [am__tar='$${TAR-tar} chof - "$$tardir"' am__untar='$${TAR-tar} xf -'],
+ [am__tar='${AMTAR} chof - "$$tardir"'; am__untar='${AMTAR} xf -'],
[m4_case([$1], [ustar],, [pax],,
[m4_fatal([Unknown tar format])])
AC_MSG_CHECKING([how to create a $1 tar archive])
/* Define to 1 if you have the <dlfcn.h> header file. */
#undef HAVE_DLFCN_H
+/* Define to 1 if you have the `getegid' function. */
+#undef HAVE_GETEGID
+
+/* Define to 1 if you have the `geteuid' function. */
+#undef HAVE_GETEUID
+
+/* Define to 1 if you have the `getgid' function. */
+#undef HAVE_GETGID
+
/* Define to 1 if you have the `getloadavg' function. */
#undef HAVE_GETLOADAVG
+/* Define to 1 if you have the `getuid' function. */
+#undef HAVE_GETUID
+
/* Define to 1 if you have the <inttypes.h> header file. */
#undef HAVE_INTTYPES_H
/* Define to 1 if you have the <pthread.h> header file. */
#undef HAVE_PTHREAD_H
+/* Define to 1 if you have the `secure_getenv' function. */
+#undef HAVE_SECURE_GETENV
+
/* Define to 1 if you have the <semaphore.h> header file. */
#undef HAVE_SEMAPHORE_H
/* Define to 1 if you have the <unistd.h> header file. */
#undef HAVE_UNISTD_H
+/* Define to 1 if you have the `__secure_getenv' function. */
+#undef HAVE___SECURE_GETENV
+
+/* Define path to HSA runtime. */
+#undef HSA_RUNTIME_LIB
+
/* Define to 1 if GNU symbol versioning is used for libgomp. */
#undef LIBGOMP_GNU_SYMBOL_VERSIONING
# include <unistd.h>
#endif"
+ac_header_list=
+ac_func_list=
ac_subst_vars='am__EXEEXT_FALSE
am__EXEEXT_TRUE
LTLIBOBJS
PLUGIN_HSA_LDFLAGS
PLUGIN_HSA_CPPFLAGS
PLUGIN_HSA
-HSA_KMT_LIB
HSA_RUNTIME_LIB
HSA_RUNTIME_INCLUDE
PLUGIN_NVPTX_LIBS
am__fastdepCC_FALSE
am__fastdepCC_TRUE
CCDEPMODE
-am__nodep
AMDEPBACKSLASH
AMDEP_FALSE
AMDEP_TRUE
with_hsa_runtime
with_hsa_runtime_include
with_hsa_runtime_lib
-with_hsa_kmt_lib
enable_linux_futex
enable_tls
enable_symvers
--with-hsa-runtime-lib=PATH
specify directory for the installed HSA run-time
library
- --with-hsa-kmt-lib=PATH specify directory for installed HSA KMT library.
Some influential environment variables:
CC C compiler command
>$cache_file
fi
+as_fn_append ac_header_list " unistd.h"
+as_fn_append ac_func_list " secure_getenv"
+as_fn_append ac_func_list " __secure_getenv"
+as_fn_append ac_func_list " getuid"
+as_fn_append ac_func_list " geteuid"
+as_fn_append ac_func_list " getgid"
+as_fn_append ac_func_list " getegid"
# Check that the precious variables saved in the cache have kept the same
# value.
ac_cache_corrupted=false
# We need awk for the "check" target. The system "awk" is bad on
# some platforms.
-# Always define AMTAR for backward compatibility. Yes, it's still used
-# in the wild :-( We should find a proper way to deprecate it ...
-AMTAR='$${TAR-tar}'
+# Always define AMTAR for backward compatibility.
-am__tar='$${TAR-tar} chof - "$$tardir"' am__untar='$${TAR-tar} xf -'
+AMTAR=${AMTAR-"${am_missing_run}tar"}
+
+am__tar='${AMTAR} chof - "$$tardir"'; am__untar='${AMTAR} xf -'
if test "x$enable_dependency_tracking" != xno; then
am_depcomp="$ac_aux_dir/depcomp"
AMDEPBACKSLASH='\'
- am__nodep='_no'
fi
if test "x$enable_dependency_tracking" != xno; then
AMDEP_TRUE=
# instance it was reported that on HP-UX the gcc test will end up
# making a dummy file named `D' -- because `-MD' means `put the output
# in D'.
- rm -rf conftest.dir
mkdir conftest.dir
# Copy depcomp to subdir because otherwise we won't find it if we're
# using a relative directory.
break
fi
;;
- msvc7 | msvc7msys | msvisualcpp | msvcmsys)
+ msvisualcpp | msvcmsys)
# This compiler won't grok `-c -o', but also, the minuso test has
# not run yet. These depmodes are late enough in the game, and
# so weak that their functioning should not be impacted.
lt_dlunknown=0; lt_dlno_uscore=1; lt_dlneed_uscore=2
lt_status=$lt_dlunknown
cat > conftest.$ac_ext <<_LT_EOF
-#line 11148 "configure"
+#line 11151 "configure"
#include "confdefs.h"
#if HAVE_DLFCN_H
lt_dlunknown=0; lt_dlno_uscore=1; lt_dlneed_uscore=2
lt_status=$lt_dlunknown
cat > conftest.$ac_ext <<_LT_EOF
-#line 11254 "configure"
+#line 11257 "configure"
#include "confdefs.h"
#if HAVE_DLFCN_H
as_fn_error "Can't support offloading without support for plugins" "$LINENO" 5
fi
+
+
+ for ac_header in $ac_header_list
+do :
+ as_ac_Header=`$as_echo "ac_cv_header_$ac_header" | $as_tr_sh`
+ac_fn_c_check_header_compile "$LINENO" "$ac_header" "$as_ac_Header" "$ac_includes_default
+"
+eval as_val=\$$as_ac_Header
+ if test "x$as_val" = x""yes; then :
+ cat >>confdefs.h <<_ACEOF
+#define `$as_echo "HAVE_$ac_header" | $as_tr_cpp` 1
+_ACEOF
+
+fi
+
+done
+
+
+
+
+
+
+
+ for ac_func in $ac_func_list
+do :
+ as_ac_var=`$as_echo "ac_cv_func_$ac_func" | $as_tr_sh`
+ac_fn_c_check_func "$LINENO" "$ac_func" "$as_ac_var"
+eval as_val=\$$as_ac_var
+ if test "x$as_val" = x""yes; then :
+ cat >>confdefs.h <<_ACEOF
+#define `$as_echo "HAVE_$ac_func" | $as_tr_cpp` 1
+_ACEOF
+
+fi
+done
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
# Look for the CUDA driver package.
CUDA_DRIVER_INCLUDE=
CUDA_DRIVER_LIB=
HSA_RUNTIME_LDFLAGS=-L$HSA_RUNTIME_LIB
fi
-HSA_KMT_LIB=
-
-HSA_KMT_LDFLAGS=
-
-# Check whether --with-hsa-kmt-lib was given.
-if test "${with_hsa_kmt_lib+set}" = set; then :
- withval=$with_hsa_kmt_lib;
-fi
-
-if test "x$with_hsa_kmt_lib" != x; then
- HSA_KMT_LIB=$with_hsa_kmt_lib
-fi
-if test "x$HSA_KMT_LIB" != x; then
- HSA_KMT_LDFLAGS=-L$HSA_KMT_LIB
-fi
-
PLUGIN_HSA=0
PLUGIN_HSA_CPPFLAGS=
PLUGIN_HSA_LDFLAGS=
-
-
# Get offload targets and path to install tree of offloading compiler.
offload_additional_options=
offload_additional_lib_paths=
tgt_name=hsa
PLUGIN_HSA=$tgt
PLUGIN_HSA_CPPFLAGS=$HSA_RUNTIME_CPPFLAGS
- PLUGIN_HSA_LDFLAGS="$HSA_RUNTIME_LDFLAGS $HSA_KMT_LDFLAGS"
- PLUGIN_HSA_LIBS="-lhsa-runtime64 -lhsakmt"
+ PLUGIN_HSA_LDFLAGS="$HSA_RUNTIME_LDFLAGS"
+ PLUGIN_HSA_LIBS="-ldl"
PLUGIN_HSA_save_CPPFLAGS=$CPPFLAGS
CPPFLAGS="$PLUGIN_HSA_CPPFLAGS $CPPFLAGS"
PLUGIN_HSA_save_LIBS=$LIBS
LIBS="$PLUGIN_HSA_LIBS $LIBS"
- cat confdefs.h - <<_ACEOF >conftest.$ac_ext
-/* end confdefs.h. */
-#include "hsa.h"
-int
-main ()
-{
-hsa_status_t status = hsa_init ()
- ;
- return 0;
-}
-_ACEOF
-if ac_fn_c_try_link "$LINENO"; then :
- PLUGIN_HSA=1
-fi
-rm -f core conftest.err conftest.$ac_objext \
- conftest$ac_exeext conftest.$ac_ext
+ PLUGIN_HSA=1
CPPFLAGS=$PLUGIN_HSA_save_CPPFLAGS
LDFLAGS=$PLUGIN_HSA_save_LDFLAGS
LIBS=$PLUGIN_HSA_save_LIBS
_ACEOF
+if test "$HSA_RUNTIME_LIB" != ""; then
+ HSA_RUNTIME_LIB="$HSA_RUNTIME_LIB/"
+fi
+
+
+cat >>confdefs.h <<_ACEOF
+#define HSA_RUNTIME_LIB "$HSA_RUNTIME_LIB"
+_ACEOF
+
+
# Check for functions needed.
for ac_func in getloadavg clock_gettime strtoull
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)
+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)
elif test "x${enable_offload_targets-no}" != xno; then
AC_MSG_ERROR([Can't support offloading without support for plugins])
fi
+AC_CHECK_HEADERS_ONCE(unistd.h)
+AC_CHECK_FUNCS_ONCE(secure_getenv __secure_getenv getuid geteuid getgid getegid)
+
# Look for the CUDA driver package.
CUDA_DRIVER_INCLUDE=
HSA_RUNTIME_LDFLAGS=-L$HSA_RUNTIME_LIB
fi
-HSA_KMT_LIB=
-AC_SUBST(HSA_KMT_LIB)
-HSA_KMT_LDFLAGS=
-AC_ARG_WITH(hsa-kmt-lib,
- [AS_HELP_STRING([--with-hsa-kmt-lib=PATH],
- [specify directory for installed HSA KMT library.])])
-if test "x$with_hsa_kmt_lib" != x; then
- HSA_KMT_LIB=$with_hsa_kmt_lib
-fi
-if test "x$HSA_KMT_LIB" != x; then
- HSA_KMT_LDFLAGS=-L$HSA_KMT_LIB
-fi
-
PLUGIN_HSA=0
PLUGIN_HSA_CPPFLAGS=
PLUGIN_HSA_LDFLAGS=
AC_SUBST(PLUGIN_HSA_LDFLAGS)
AC_SUBST(PLUGIN_HSA_LIBS)
-
-
# Get offload targets and path to install tree of offloading compiler.
offload_additional_options=
offload_additional_lib_paths=
tgt_name=hsa
PLUGIN_HSA=$tgt
PLUGIN_HSA_CPPFLAGS=$HSA_RUNTIME_CPPFLAGS
- PLUGIN_HSA_LDFLAGS="$HSA_RUNTIME_LDFLAGS $HSA_KMT_LDFLAGS"
- PLUGIN_HSA_LIBS="-lhsa-runtime64 -lhsakmt"
+ PLUGIN_HSA_LDFLAGS="$HSA_RUNTIME_LDFLAGS"
+ PLUGIN_HSA_LIBS="-ldl"
PLUGIN_HSA_save_CPPFLAGS=$CPPFLAGS
CPPFLAGS="$PLUGIN_HSA_CPPFLAGS $CPPFLAGS"
PLUGIN_HSA_save_LIBS=$LIBS
LIBS="$PLUGIN_HSA_LIBS $LIBS"
- AC_LINK_IFELSE(
- [AC_LANG_PROGRAM(
- [#include "hsa.h"],
- [hsa_status_t status = hsa_init ()])],
- [PLUGIN_HSA=1])
+ PLUGIN_HSA=1
CPPFLAGS=$PLUGIN_HSA_save_CPPFLAGS
LDFLAGS=$PLUGIN_HSA_save_LDFLAGS
LIBS=$PLUGIN_HSA_save_LIBS
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.])
+
+if test "$HSA_RUNTIME_LIB" != ""; then
+ HSA_RUNTIME_LIB="$HSA_RUNTIME_LIB/"
+fi
+
+AC_DEFINE_UNQUOTED([HSA_RUNTIME_LIB], ["$HSA_RUNTIME_LIB"],
+ [Define path to HSA runtime.])
--- /dev/null
+/* HSA runtime API 1.0.1 representation description.
+ Copyright (C) 2016 Free Software Foundation, Inc.
+
+This file is part of GCC.
+
+GCC is free software; you can redistribute it and/or modify
+it under the terms of the GNU General Public License as published by
+the Free Software Foundation; either version 3, or (at your option)
+any later version.
+
+GCC is distributed in the hope that it will be useful,
+but WITHOUT ANY WARRANTY; without even the implied warranty of
+MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
+GNU General Public License for more details.
+
+You should have received a copy of the GNU General Public License
+along with GCC; see the file COPYING3. If not see
+<http://www.gnu.org/licenses/>.
+
+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_H
+#define _HSA_H 1
+
+#define HSA_LARGE_MODEL 1
+
+typedef struct hsa_signal_s { uint64_t handle; } hsa_signal_t;
+typedef enum {
+ HSA_QUEUE_TYPE_MULTI = 0,
+ HSA_QUEUE_TYPE_SINGLE = 1
+} hsa_queue_type_t;
+
+typedef enum { HSA_PROFILE_BASE = 0, HSA_PROFILE_FULL = 1 } hsa_profile_t;
+typedef struct hsa_region_s { uint64_t handle; } hsa_region_t;
+typedef enum {
+ HSA_EXECUTABLE_SYMBOL_INFO_TYPE = 0,
+ HSA_EXECUTABLE_SYMBOL_INFO_NAME_LENGTH = 1,
+ HSA_EXECUTABLE_SYMBOL_INFO_NAME = 2,
+ HSA_EXECUTABLE_SYMBOL_INFO_MODULE_NAME_LENGTH = 3,
+ HSA_EXECUTABLE_SYMBOL_INFO_MODULE_NAME = 4,
+ HSA_EXECUTABLE_SYMBOL_INFO_AGENT = 20,
+ HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_ADDRESS = 21,
+ HSA_EXECUTABLE_SYMBOL_INFO_LINKAGE = 5,
+ HSA_EXECUTABLE_SYMBOL_INFO_IS_DEFINITION = 17,
+ HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_ALLOCATION = 6,
+ HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_SEGMENT = 7,
+ HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_ALIGNMENT = 8,
+ HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_SIZE = 9,
+ HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_IS_CONST = 10,
+ HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT = 22,
+ HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_KERNARG_SEGMENT_SIZE = 11,
+ HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_KERNARG_SEGMENT_ALIGNMENT = 12,
+ HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_GROUP_SEGMENT_SIZE = 13,
+ HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_PRIVATE_SEGMENT_SIZE = 14,
+ HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_DYNAMIC_CALLSTACK = 15,
+ HSA_EXECUTABLE_SYMBOL_INFO_INDIRECT_FUNCTION_OBJECT = 23,
+ HSA_EXECUTABLE_SYMBOL_INFO_INDIRECT_FUNCTION_CALL_CONVENTION = 16
+} hsa_executable_symbol_info_t;
+typedef enum {
+ HSA_REGION_GLOBAL_FLAG_KERNARG = 1,
+ HSA_REGION_GLOBAL_FLAG_FINE_GRAINED = 2,
+ HSA_REGION_GLOBAL_FLAG_COARSE_GRAINED = 4
+} hsa_region_global_flag_t;
+typedef struct hsa_code_object_s { uint64_t handle; } hsa_code_object_t;
+typedef enum {
+ HSA_KERNEL_DISPATCH_PACKET_SETUP_WIDTH_DIMENSIONS = 2
+} hsa_kernel_dispatch_packet_setup_width_t;
+typedef enum {
+ HSA_DEVICE_TYPE_CPU = 0,
+ HSA_DEVICE_TYPE_GPU = 1,
+ HSA_DEVICE_TYPE_DSP = 2
+} hsa_device_type_t;
+typedef enum {
+ HSA_STATUS_SUCCESS = 0x0,
+ HSA_STATUS_INFO_BREAK = 0x1,
+ HSA_STATUS_ERROR = 0x1000,
+ HSA_STATUS_ERROR_INVALID_ARGUMENT = 0x1001,
+ HSA_STATUS_ERROR_INVALID_QUEUE_CREATION = 0x1002,
+ HSA_STATUS_ERROR_INVALID_ALLOCATION = 0x1003,
+ HSA_STATUS_ERROR_INVALID_AGENT = 0x1004,
+ HSA_STATUS_ERROR_INVALID_REGION = 0x1005,
+ HSA_STATUS_ERROR_INVALID_SIGNAL = 0x1006,
+ HSA_STATUS_ERROR_INVALID_QUEUE = 0x1007,
+ HSA_STATUS_ERROR_OUT_OF_RESOURCES = 0x1008,
+ HSA_STATUS_ERROR_INVALID_PACKET_FORMAT = 0x1009,
+ HSA_STATUS_ERROR_RESOURCE_FREE = 0x100A,
+ HSA_STATUS_ERROR_NOT_INITIALIZED = 0x100B,
+ HSA_STATUS_ERROR_REFCOUNT_OVERFLOW = 0x100C,
+ HSA_STATUS_ERROR_INCOMPATIBLE_ARGUMENTS = 0x100D,
+ HSA_STATUS_ERROR_INVALID_INDEX = 0x100E,
+ HSA_STATUS_ERROR_INVALID_ISA = 0x100F,
+ HSA_STATUS_ERROR_INVALID_ISA_NAME = 0x1017,
+ HSA_STATUS_ERROR_INVALID_CODE_OBJECT = 0x1010,
+ HSA_STATUS_ERROR_INVALID_EXECUTABLE = 0x1011,
+ HSA_STATUS_ERROR_FROZEN_EXECUTABLE = 0x1012,
+ HSA_STATUS_ERROR_INVALID_SYMBOL_NAME = 0x1013,
+ HSA_STATUS_ERROR_VARIABLE_ALREADY_DEFINED = 0x1014,
+ HSA_STATUS_ERROR_VARIABLE_UNDEFINED = 0x1015,
+ HSA_STATUS_ERROR_EXCEPTION = 0x1016
+} hsa_status_t;
+typedef enum {
+ HSA_EXTENSION_FINALIZER = 0,
+ HSA_EXTENSION_IMAGES = 1
+} hsa_extension_t;
+typedef struct hsa_queue_s {
+ hsa_queue_type_t type;
+ uint32_t features;
+
+#ifdef HSA_LARGE_MODEL
+ void *base_address;
+#elif defined HSA_LITTLE_ENDIAN
+ void *base_address;
+ uint32_t reserved0;
+#else
+ uint32_t reserved0;
+ void *base_address;
+#endif
+
+ hsa_signal_t doorbell_signal;
+ uint32_t size;
+ uint32_t reserved1;
+ uint64_t id;
+} hsa_queue_t;
+typedef struct hsa_agent_dispatch_packet_s {
+ uint16_t header;
+ uint16_t type;
+ uint32_t reserved0;
+
+#ifdef HSA_LARGE_MODEL
+ void *return_address;
+#elif defined HSA_LITTLE_ENDIAN
+ void *return_address;
+ uint32_t reserved1;
+#else
+ uint32_t reserved1;
+ void *return_address;
+#endif
+ uint64_t arg[4];
+ uint64_t reserved2;
+ hsa_signal_t completion_signal;
+} hsa_agent_dispatch_packet_t;
+typedef enum {
+ HSA_CODE_SYMBOL_INFO_TYPE = 0,
+ HSA_CODE_SYMBOL_INFO_NAME_LENGTH = 1,
+ HSA_CODE_SYMBOL_INFO_NAME = 2,
+ HSA_CODE_SYMBOL_INFO_MODULE_NAME_LENGTH = 3,
+ HSA_CODE_SYMBOL_INFO_MODULE_NAME = 4,
+ HSA_CODE_SYMBOL_INFO_LINKAGE = 5,
+ HSA_CODE_SYMBOL_INFO_IS_DEFINITION = 17,
+ HSA_CODE_SYMBOL_INFO_VARIABLE_ALLOCATION = 6,
+ HSA_CODE_SYMBOL_INFO_VARIABLE_SEGMENT = 7,
+ HSA_CODE_SYMBOL_INFO_VARIABLE_ALIGNMENT = 8,
+ HSA_CODE_SYMBOL_INFO_VARIABLE_SIZE = 9,
+ HSA_CODE_SYMBOL_INFO_VARIABLE_IS_CONST = 10,
+ HSA_CODE_SYMBOL_INFO_KERNEL_KERNARG_SEGMENT_SIZE = 11,
+ HSA_CODE_SYMBOL_INFO_KERNEL_KERNARG_SEGMENT_ALIGNMENT = 12,
+ HSA_CODE_SYMBOL_INFO_KERNEL_GROUP_SEGMENT_SIZE = 13,
+ HSA_CODE_SYMBOL_INFO_KERNEL_PRIVATE_SEGMENT_SIZE = 14,
+ HSA_CODE_SYMBOL_INFO_KERNEL_DYNAMIC_CALLSTACK = 15,
+ HSA_CODE_SYMBOL_INFO_INDIRECT_FUNCTION_CALL_CONVENTION = 16
+} hsa_code_symbol_info_t;
+typedef enum {
+ HSA_QUEUE_FEATURE_KERNEL_DISPATCH = 1,
+ HSA_QUEUE_FEATURE_AGENT_DISPATCH = 2
+} hsa_queue_feature_t;
+typedef enum {
+ HSA_VARIABLE_ALLOCATION_AGENT = 0,
+ HSA_VARIABLE_ALLOCATION_PROGRAM = 1
+} hsa_variable_allocation_t;
+typedef enum {
+ HSA_FENCE_SCOPE_NONE = 0,
+ HSA_FENCE_SCOPE_AGENT = 1,
+ HSA_FENCE_SCOPE_SYSTEM = 2
+} hsa_fence_scope_t;
+typedef struct hsa_agent_s { uint64_t handle; } hsa_agent_t;
+typedef enum { HSA_CODE_OBJECT_TYPE_PROGRAM = 0 } hsa_code_object_type_t;
+typedef enum {
+ HSA_SIGNAL_CONDITION_EQ = 0,
+ HSA_SIGNAL_CONDITION_NE = 1,
+ HSA_SIGNAL_CONDITION_LT = 2,
+ HSA_SIGNAL_CONDITION_GTE = 3
+} hsa_signal_condition_t;
+typedef enum {
+ HSA_EXECUTABLE_STATE_UNFROZEN = 0,
+ HSA_EXECUTABLE_STATE_FROZEN = 1
+} hsa_executable_state_t;
+typedef enum {
+ HSA_ENDIANNESS_LITTLE = 0,
+ HSA_ENDIANNESS_BIG = 1
+} hsa_endianness_t;
+typedef enum {
+ HSA_MACHINE_MODEL_SMALL = 0,
+ HSA_MACHINE_MODEL_LARGE = 1
+} hsa_machine_model_t;
+typedef enum {
+ HSA_AGENT_INFO_NAME = 0,
+ HSA_AGENT_INFO_VENDOR_NAME = 1,
+ HSA_AGENT_INFO_FEATURE = 2,
+ HSA_AGENT_INFO_MACHINE_MODEL = 3,
+ HSA_AGENT_INFO_PROFILE = 4,
+ HSA_AGENT_INFO_DEFAULT_FLOAT_ROUNDING_MODE = 5,
+ HSA_AGENT_INFO_BASE_PROFILE_DEFAULT_FLOAT_ROUNDING_MODES = 23,
+ HSA_AGENT_INFO_FAST_F16_OPERATION = 24,
+ HSA_AGENT_INFO_WAVEFRONT_SIZE = 6,
+ HSA_AGENT_INFO_WORKGROUP_MAX_DIM = 7,
+ HSA_AGENT_INFO_WORKGROUP_MAX_SIZE = 8,
+ HSA_AGENT_INFO_GRID_MAX_DIM = 9,
+ HSA_AGENT_INFO_GRID_MAX_SIZE = 10,
+ HSA_AGENT_INFO_FBARRIER_MAX_SIZE = 11,
+ HSA_AGENT_INFO_QUEUES_MAX = 12,
+ HSA_AGENT_INFO_QUEUE_MIN_SIZE = 13,
+ HSA_AGENT_INFO_QUEUE_MAX_SIZE = 14,
+ HSA_AGENT_INFO_QUEUE_TYPE = 15,
+ HSA_AGENT_INFO_NODE = 16,
+ HSA_AGENT_INFO_DEVICE = 17,
+ HSA_AGENT_INFO_CACHE_SIZE = 18,
+ HSA_AGENT_INFO_ISA = 19,
+ HSA_AGENT_INFO_EXTENSIONS = 20,
+ HSA_AGENT_INFO_VERSION_MAJOR = 21,
+ HSA_AGENT_INFO_VERSION_MINOR = 22
+} hsa_agent_info_t;
+typedef struct hsa_barrier_and_packet_s {
+ uint16_t header;
+ uint16_t reserved0;
+ uint32_t reserved1;
+ hsa_signal_t dep_signal[5];
+ uint64_t reserved2;
+ hsa_signal_t completion_signal;
+} hsa_barrier_and_packet_t;
+typedef struct hsa_dim3_s {
+ uint32_t x;
+ uint32_t y;
+ uint32_t z;
+} hsa_dim3_t;
+typedef enum {
+ HSA_ACCESS_PERMISSION_RO = 1,
+ HSA_ACCESS_PERMISSION_WO = 2,
+ HSA_ACCESS_PERMISSION_RW = 3
+} hsa_access_permission_t;
+typedef enum {
+ HSA_AGENT_FEATURE_KERNEL_DISPATCH = 1,
+ HSA_AGENT_FEATURE_AGENT_DISPATCH = 2
+} hsa_agent_feature_t;
+typedef enum {
+ HSA_WAIT_STATE_BLOCKED = 0,
+ HSA_WAIT_STATE_ACTIVE = 1
+} hsa_wait_state_t;
+typedef struct hsa_executable_s { uint64_t handle; } hsa_executable_t;
+typedef enum {
+ HSA_REGION_SEGMENT_GLOBAL = 0,
+ HSA_REGION_SEGMENT_READONLY = 1,
+ HSA_REGION_SEGMENT_PRIVATE = 2,
+ HSA_REGION_SEGMENT_GROUP = 3
+} hsa_region_segment_t;
+typedef enum {
+ HSA_REGION_INFO_SEGMENT = 0,
+ HSA_REGION_INFO_GLOBAL_FLAGS = 1,
+ HSA_REGION_INFO_SIZE = 2,
+ HSA_REGION_INFO_ALLOC_MAX_SIZE = 4,
+ HSA_REGION_INFO_RUNTIME_ALLOC_ALLOWED = 5,
+ HSA_REGION_INFO_RUNTIME_ALLOC_GRANULE = 6,
+ HSA_REGION_INFO_RUNTIME_ALLOC_ALIGNMENT = 7
+} hsa_region_info_t;
+typedef enum {
+ HSA_ISA_INFO_NAME_LENGTH = 0,
+ HSA_ISA_INFO_NAME = 1,
+ HSA_ISA_INFO_CALL_CONVENTION_COUNT = 2,
+ HSA_ISA_INFO_CALL_CONVENTION_INFO_WAVEFRONT_SIZE = 3,
+ HSA_ISA_INFO_CALL_CONVENTION_INFO_WAVEFRONTS_PER_COMPUTE_UNIT = 4
+} hsa_isa_info_t;
+typedef enum {
+ HSA_VARIABLE_SEGMENT_GLOBAL = 0,
+ HSA_VARIABLE_SEGMENT_READONLY = 1
+} hsa_variable_segment_t;
+typedef struct hsa_callback_data_s { uint64_t handle; } hsa_callback_data_t;
+typedef enum {
+ HSA_SYMBOL_KIND_VARIABLE = 0,
+ HSA_SYMBOL_KIND_KERNEL = 1,
+ HSA_SYMBOL_KIND_INDIRECT_FUNCTION = 2
+} hsa_symbol_kind_t;
+typedef struct hsa_kernel_dispatch_packet_s {
+ uint16_t header;
+ uint16_t setup;
+ uint16_t workgroup_size_x;
+ uint16_t workgroup_size_y;
+ uint16_t workgroup_size_z;
+ uint16_t reserved0;
+ uint32_t grid_size_x;
+ uint32_t grid_size_y;
+ uint32_t grid_size_z;
+ uint32_t private_segment_size;
+ uint32_t group_segment_size;
+ uint64_t kernel_object;
+
+#ifdef HSA_LARGE_MODEL
+ void *kernarg_address;
+#elif defined HSA_LITTLE_ENDIAN
+ void *kernarg_address;
+ uint32_t reserved1;
+#else
+ uint32_t reserved1;
+ void *kernarg_address;
+#endif
+ uint64_t reserved2;
+ hsa_signal_t completion_signal;
+} hsa_kernel_dispatch_packet_t;
+typedef enum {
+ HSA_PACKET_TYPE_VENDOR_SPECIFIC = 0,
+ HSA_PACKET_TYPE_INVALID = 1,
+ HSA_PACKET_TYPE_KERNEL_DISPATCH = 2,
+ HSA_PACKET_TYPE_BARRIER_AND = 3,
+ HSA_PACKET_TYPE_AGENT_DISPATCH = 4,
+ HSA_PACKET_TYPE_BARRIER_OR = 5
+} hsa_packet_type_t;
+typedef enum {
+ HSA_PACKET_HEADER_TYPE = 0,
+ HSA_PACKET_HEADER_BARRIER = 8,
+ HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE = 9,
+ HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE = 11
+} hsa_packet_header_t;
+typedef struct hsa_isa_s { uint64_t handle; } hsa_isa_t;
+typedef enum {
+ HSA_DEFAULT_FLOAT_ROUNDING_MODE_DEFAULT = 0,
+ HSA_DEFAULT_FLOAT_ROUNDING_MODE_ZERO = 1,
+ HSA_DEFAULT_FLOAT_ROUNDING_MODE_NEAR = 2
+} hsa_default_float_rounding_mode_t;
+typedef struct hsa_code_symbol_s { uint64_t handle; } hsa_code_symbol_t;
+typedef struct hsa_executable_symbol_s {
+ uint64_t handle;
+} hsa_executable_symbol_t;
+#ifdef HSA_LARGE_MODEL
+typedef int64_t hsa_signal_value_t;
+#else
+typedef int32_t hsa_signal_value_t;
+#endif
+typedef enum {
+ HSA_EXCEPTION_POLICY_BREAK = 1,
+ HSA_EXCEPTION_POLICY_DETECT = 2
+} hsa_exception_policy_t;
+typedef enum {
+ HSA_SYSTEM_INFO_VERSION_MAJOR = 0,
+ HSA_SYSTEM_INFO_VERSION_MINOR = 1,
+ HSA_SYSTEM_INFO_TIMESTAMP = 2,
+ HSA_SYSTEM_INFO_TIMESTAMP_FREQUENCY = 3,
+ HSA_SYSTEM_INFO_SIGNAL_MAX_WAIT = 4,
+ HSA_SYSTEM_INFO_ENDIANNESS = 5,
+ HSA_SYSTEM_INFO_MACHINE_MODEL = 6,
+ HSA_SYSTEM_INFO_EXTENSIONS = 7
+} hsa_system_info_t;
+typedef enum {
+ HSA_EXECUTABLE_INFO_PROFILE = 1,
+ HSA_EXECUTABLE_INFO_STATE = 2
+} hsa_executable_info_t;
+typedef enum {
+ HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS = 0
+} hsa_kernel_dispatch_packet_setup_t;
+typedef enum {
+ HSA_PACKET_HEADER_WIDTH_TYPE = 8,
+ HSA_PACKET_HEADER_WIDTH_BARRIER = 1,
+ HSA_PACKET_HEADER_WIDTH_ACQUIRE_FENCE_SCOPE = 2,
+ HSA_PACKET_HEADER_WIDTH_RELEASE_FENCE_SCOPE = 2
+} hsa_packet_header_width_t;
+typedef enum {
+ HSA_CODE_OBJECT_INFO_VERSION = 0,
+ HSA_CODE_OBJECT_INFO_TYPE = 1,
+ HSA_CODE_OBJECT_INFO_ISA = 2,
+ HSA_CODE_OBJECT_INFO_MACHINE_MODEL = 3,
+ HSA_CODE_OBJECT_INFO_PROFILE = 4,
+ HSA_CODE_OBJECT_INFO_DEFAULT_FLOAT_ROUNDING_MODE = 5
+} hsa_code_object_info_t;
+typedef struct hsa_barrier_or_packet_s {
+ uint16_t header;
+ uint16_t reserved0;
+ uint32_t reserved1;
+ hsa_signal_t dep_signal[5];
+ uint64_t reserved2;
+ hsa_signal_t completion_signal;
+} hsa_barrier_or_packet_t;
+typedef enum {
+ HSA_SYMBOL_KIND_LINKAGE_MODULE = 0,
+ HSA_SYMBOL_KIND_LINKAGE_PROGRAM = 1,
+} hsa_symbol_kind_linkage_t;
+hsa_status_t hsa_executable_validate(hsa_executable_t executable,
+ uint32_t *result);
+uint64_t hsa_queue_add_write_index_acq_rel(const hsa_queue_t *queue,
+ uint64_t value);
+
+uint64_t hsa_queue_add_write_index_acquire(const hsa_queue_t *queue,
+ uint64_t value);
+
+uint64_t hsa_queue_add_write_index_relaxed(const hsa_queue_t *queue,
+ uint64_t value);
+
+uint64_t hsa_queue_add_write_index_release(const hsa_queue_t *queue,
+ uint64_t value);
+hsa_status_t hsa_shut_down();
+void hsa_signal_add_acq_rel(hsa_signal_t signal, hsa_signal_value_t value);
+
+void hsa_signal_add_acquire(hsa_signal_t signal, hsa_signal_value_t value);
+
+void hsa_signal_add_relaxed(hsa_signal_t signal, hsa_signal_value_t value);
+
+void hsa_signal_add_release(hsa_signal_t signal, hsa_signal_value_t value);
+hsa_status_t hsa_executable_readonly_variable_define(
+ hsa_executable_t executable, hsa_agent_t agent, const char *variable_name,
+ void *address);
+hsa_status_t hsa_agent_extension_supported(uint16_t extension,
+ hsa_agent_t agent,
+ uint16_t version_major,
+ uint16_t version_minor,
+ bool *result);
+hsa_signal_value_t hsa_signal_load_acquire(hsa_signal_t signal);
+
+hsa_signal_value_t hsa_signal_load_relaxed(hsa_signal_t signal);
+hsa_status_t hsa_executable_get_info(hsa_executable_t executable,
+ hsa_executable_info_t attribute,
+ void *value);
+hsa_status_t hsa_iterate_agents(hsa_status_t (*callback)(hsa_agent_t agent,
+ void *data),
+ void *data);
+void hsa_signal_subtract_acq_rel(hsa_signal_t signal, hsa_signal_value_t value);
+
+void hsa_signal_subtract_acquire(hsa_signal_t signal, hsa_signal_value_t value);
+
+void hsa_signal_subtract_relaxed(hsa_signal_t signal, hsa_signal_value_t value);
+
+void hsa_signal_subtract_release(hsa_signal_t signal, hsa_signal_value_t value);
+hsa_status_t
+hsa_executable_symbol_get_info(hsa_executable_symbol_t executable_symbol,
+ hsa_executable_symbol_info_t attribute,
+ void *value);
+void hsa_signal_xor_acq_rel(hsa_signal_t signal, hsa_signal_value_t value);
+
+void hsa_signal_xor_acquire(hsa_signal_t signal, hsa_signal_value_t value);
+
+void hsa_signal_xor_relaxed(hsa_signal_t signal, hsa_signal_value_t value);
+
+void hsa_signal_xor_release(hsa_signal_t signal, hsa_signal_value_t value);
+hsa_status_t hsa_code_object_get_info(hsa_code_object_t code_object,
+ hsa_code_object_info_t attribute,
+ void *value);
+hsa_status_t hsa_code_object_deserialize(void *serialized_code_object,
+ size_t serialized_code_object_size,
+ const char *options,
+ hsa_code_object_t *code_object);
+hsa_status_t hsa_status_string(hsa_status_t status, const char **status_string);
+hsa_status_t hsa_code_object_get_symbol(hsa_code_object_t code_object,
+ const char *symbol_name,
+ hsa_code_symbol_t *symbol);
+void hsa_signal_store_relaxed(hsa_signal_t signal, hsa_signal_value_t value);
+
+void hsa_signal_store_release(hsa_signal_t signal, hsa_signal_value_t value);
+hsa_status_t hsa_signal_destroy(hsa_signal_t signal);
+hsa_status_t hsa_system_get_extension_table(uint16_t extension,
+ uint16_t version_major,
+ uint16_t version_minor,
+ void *table);
+hsa_status_t hsa_agent_iterate_regions(
+ hsa_agent_t agent,
+ hsa_status_t (*callback)(hsa_region_t region, void *data), void *data);
+hsa_status_t hsa_executable_agent_global_variable_define(
+ hsa_executable_t executable, hsa_agent_t agent, const char *variable_name,
+ void *address);
+hsa_status_t hsa_queue_create(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_isa_compatible(hsa_isa_t code_object_isa, hsa_isa_t agent_isa,
+ bool *result);
+hsa_status_t hsa_code_object_serialize(
+ hsa_code_object_t code_object,
+ hsa_status_t (*alloc_callback)(size_t size, hsa_callback_data_t data,
+ void **address),
+ hsa_callback_data_t callback_data, const char *options,
+ void **serialized_code_object, size_t *serialized_code_object_size);
+hsa_status_t hsa_region_get_info(hsa_region_t region,
+ hsa_region_info_t attribute, void *value);
+hsa_status_t hsa_executable_freeze(hsa_extension_t executable,
+ const char *options);
+hsa_status_t hsa_system_extension_supported(uint16_t extension,
+ uint16_t version_major,
+ uint16_t version_minor,
+ bool *result);
+hsa_signal_value_t hsa_signal_wait_acquire(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_wait_relaxed(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_status_t hsa_memory_copy(void *dst, const void *src, size_t size);
+hsa_status_t hsa_memory_free(void *ptr);
+hsa_status_t hsa_queue_destroy(hsa_queue_t *queue);
+hsa_status_t hsa_isa_from_name(const char *name, hsa_isa_t *isa);
+hsa_status_t hsa_isa_get_info(hsa_isa_t isa, hsa_isa_info_t attribute,
+ uint32_t index, void *value);
+hsa_status_t hsa_signal_create(hsa_signal_value_t initial_value,
+ uint32_t num_consumers,
+ const hsa_agent_t *consumers,
+ hsa_signal_t *signal);
+hsa_status_t hsa_code_symbol_get_info(hsa_code_symbol_t code_symbol,
+ hsa_code_symbol_info_t attribute,
+ void *value);
+hsa_signal_value_t hsa_signal_cas_acq_rel(hsa_signal_t signal,
+ hsa_signal_value_t expected,
+ hsa_signal_value_t value);
+
+hsa_signal_value_t hsa_signal_cas_acquire(hsa_signal_t signal,
+ hsa_signal_value_t expected,
+ hsa_signal_value_t value);
+
+hsa_signal_value_t hsa_signal_cas_relaxed(hsa_signal_t signal,
+ hsa_signal_value_t expected,
+ hsa_signal_value_t value);
+
+hsa_signal_value_t hsa_signal_cas_release(hsa_signal_t signal,
+ hsa_signal_value_t expected,
+ hsa_signal_value_t value);
+hsa_status_t hsa_code_object_iterate_symbols(
+ hsa_code_object_t code_object,
+ hsa_status_t (*callback)(hsa_code_object_t code_object,
+ hsa_code_symbol_t symbol, void *data),
+ void *data);
+void hsa_queue_store_read_index_relaxed(const hsa_queue_t *queue,
+ uint64_t value);
+
+void hsa_queue_store_read_index_release(const hsa_queue_t *queue,
+ uint64_t value);
+hsa_status_t hsa_memory_assign_agent(void *ptr, hsa_agent_t agent,
+ hsa_access_permission_t access);
+hsa_status_t hsa_queue_inactivate(hsa_queue_t *queue);
+hsa_status_t hsa_executable_get_symbol(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);
+uint64_t hsa_queue_cas_write_index_acq_rel(const hsa_queue_t *queue,
+ uint64_t expected, uint64_t value);
+
+uint64_t hsa_queue_cas_write_index_acquire(const hsa_queue_t *queue,
+ uint64_t expected, uint64_t value);
+
+uint64_t hsa_queue_cas_write_index_relaxed(const hsa_queue_t *queue,
+ uint64_t expected, uint64_t value);
+
+uint64_t hsa_queue_cas_write_index_release(const hsa_queue_t *queue,
+ uint64_t expected, uint64_t value);
+void hsa_signal_and_acq_rel(hsa_signal_t signal, hsa_signal_value_t value);
+
+void hsa_signal_and_acquire(hsa_signal_t signal, hsa_signal_value_t value);
+
+void hsa_signal_and_relaxed(hsa_signal_t signal, hsa_signal_value_t value);
+
+void hsa_signal_and_release(hsa_signal_t signal, hsa_signal_value_t value);
+uint64_t hsa_queue_load_read_index_acquire(const hsa_queue_t *queue);
+
+uint64_t hsa_queue_load_read_index_relaxed(const hsa_queue_t *queue);
+hsa_status_t hsa_executable_load_code_object(hsa_executable_t executable,
+ hsa_agent_t agent,
+ hsa_code_object_t code_object,
+ const char *options);
+uint64_t hsa_queue_load_write_index_acquire(const hsa_queue_t *queue);
+
+uint64_t hsa_queue_load_write_index_relaxed(const hsa_queue_t *queue);
+hsa_status_t hsa_agent_get_exception_policies(hsa_agent_t agent,
+ hsa_profile_t profile,
+ uint16_t *mask);
+hsa_status_t hsa_memory_deregister(void *ptr, size_t size);
+void hsa_signal_or_acq_rel(hsa_signal_t signal, hsa_signal_value_t value);
+
+void hsa_signal_or_acquire(hsa_signal_t signal, hsa_signal_value_t value);
+
+void hsa_signal_or_relaxed(hsa_signal_t signal, hsa_signal_value_t value);
+
+void hsa_signal_or_release(hsa_signal_t signal, hsa_signal_value_t value);
+hsa_status_t hsa_soft_queue_create(hsa_region_t region, uint32_t size,
+ hsa_queue_type_t type, uint32_t features,
+ hsa_signal_t doorbell_signal,
+ hsa_queue_t **queue);
+hsa_status_t hsa_executable_iterate_symbols(
+ hsa_executable_t executable,
+ hsa_status_t (*callback)(hsa_executable_t executable,
+ hsa_executable_symbol_t symbol, void *data),
+ void *data);
+hsa_status_t hsa_memory_register(void *ptr, size_t size);
+void hsa_queue_store_write_index_relaxed(const hsa_queue_t *queue,
+ uint64_t value);
+
+void hsa_queue_store_write_index_release(const hsa_queue_t *queue,
+ uint64_t value);
+hsa_status_t hsa_executable_global_variable_define(hsa_executable_t executable,
+ const char *variable_name,
+ void *address);
+hsa_status_t hsa_executable_destroy(hsa_executable_t executable);
+hsa_status_t hsa_code_object_destroy(hsa_code_object_t code_object);
+hsa_status_t hsa_memory_allocate(hsa_region_t region, size_t size, void **ptr);
+hsa_signal_value_t hsa_signal_exchange_acq_rel(hsa_signal_t signal,
+ hsa_signal_value_t value);
+
+hsa_signal_value_t hsa_signal_exchange_acquire(hsa_signal_t signal,
+ hsa_signal_value_t value);
+
+hsa_signal_value_t hsa_signal_exchange_relaxed(hsa_signal_t signal,
+ hsa_signal_value_t value);
+
+hsa_signal_value_t hsa_signal_exchange_release(hsa_signal_t signal,
+ hsa_signal_value_t value);
+hsa_status_t hsa_agent_get_info(hsa_agent_t agent, hsa_agent_info_t attribute,
+ void *value);
+hsa_status_t hsa_init();
+hsa_status_t hsa_system_get_info(hsa_system_info_t attribute, void *value);
+hsa_status_t hsa_executable_create(hsa_profile_t profile,
+ hsa_executable_state_t executable_state,
+ const char *options,
+ hsa_executable_t *executable);
+
+#endif /* _HSA_H */
--- /dev/null
+/* HSA Extensions API 1.0.1 representation description.
+ Copyright (C) 2016 Free Software Foundation, Inc.
+
+This file is part of GCC.
+
+GCC is free software; you can redistribute it and/or modify
+it under the terms of the GNU General Public License as published by
+the Free Software Foundation; either version 3, or (at your option)
+any later version.
+
+GCC is distributed in the hope that it will be useful,
+but WITHOUT ANY WARRANTY; without even the implied warranty of
+MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
+GNU General Public License for more details.
+
+You should have received a copy of the GNU General Public License
+along with GCC; see the file COPYING3. If not see
+<http://www.gnu.org/licenses/>.
+
+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 */
see the files COPYING3 and COPYING.RUNTIME respectively. If not, see
<http://www.gnu.org/licenses/>. */
+#include "config.h"
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <pthread.h>
-#include <hsa.h>
-#include <hsa_ext_finalize.h>
+#include <inttypes.h>
+#include <stdbool.h>
+#include <plugin/hsa.h>
+#include <plugin/hsa_ext_finalize.h>
#include <dlfcn.h>
#include "libgomp-plugin.h"
#include "gomp-constants.h"
+/* 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
+
+/* 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. */
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 (getenv ("HSA_DEBUG"))
+ if (secure_getenv ("HSA_DEBUG"))
debug = true;
else
debug = false;
- if (getenv ("HSA_SUPPRESS_HOST_FALLBACK"))
+ 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
return;
const char *hsa_error_msg;
- hsa_status_string (status, &hsa_error_msg);
+ hsa_fns.hsa_status_string_fn (status, &hsa_error_msg);
fprintf (stderr, "HSA warning: %s\nRuntime message: %s", str, hsa_error_msg);
}
hsa_fatal (const char *str, hsa_status_t status)
{
const char *hsa_error_msg;
- hsa_status_string (status, &hsa_error_msg);
+ hsa_fns.hsa_status_string_fn (status, &hsa_error_msg);
GOMP_PLUGIN_fatal ("HSA fatal error: %s\nRuntime message: %s", str,
hsa_error_msg);
}
hsa_error (const char *str, hsa_status_t status)
{
const char *hsa_error_msg;
- hsa_status_string (status, &hsa_error_msg);
+ 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;
static struct hsa_context_info hsa_context;
+#define DLSYM_FN(function) \
+ hsa_fns.function##_fn = dlsym (handle, #function); \
+ if (hsa_fns.function##_fn == NULL) \
+ return false;
+
+static bool
+init_hsa_runtime_functions (void)
+{
+ void *handle = dlopen (hsa_runtime_lib, RTLD_LAZY);
+ if (handle == NULL)
+ return false;
+
+ 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;
+}
+
/* Find kernel for an AGENT by name provided in KERNEL_NAME. */
static struct kernel_info *
{
hsa_device_type_t device_type;
hsa_status_t status
- = hsa_agent_get_info (agent, HSA_AGENT_INFO_DEVICE, &device_type);
- if (status != HSA_STATUS_SUCCESS || device_type != HSA_DEVICE_TYPE_GPU)
+ = 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_agent_get_info (agent, HSA_AGENT_INFO_FEATURE, &features);
+ 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_agent_get_info (agent, HSA_AGENT_INFO_QUEUE_TYPE, &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;
if (hsa_context.initialized)
return true;
init_enviroment_variables ();
- status = hsa_init ();
+ 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_iterate_agents (count_gpu_agents, NULL);
+ 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_iterate_agents (assign_agent_ids, &agent_index);
+ 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");
hsa_status_t status;
hsa_region_segment_t segment;
- status = hsa_region_get_info (region, HSA_REGION_INFO_SEGMENT, &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_region_get_info (region, HSA_REGION_INFO_GLOBAL_FLAGS, &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)
uint32_t queue_size;
hsa_status_t status;
- status = hsa_agent_get_info (agent->id, HSA_AGENT_INFO_QUEUE_MAX_SIZE,
- &queue_size);
+ 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_agent_get_info (agent->id, HSA_AGENT_INFO_ISA, &agent->isa);
+ 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_queue_create (agent->id, queue_size, HSA_QUEUE_TYPE_MULTI,
- queue_callback, NULL, UINT32_MAX, UINT32_MAX,
- &agent->command_q);
+ 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_queue_create (agent->id, queue_size, HSA_QUEUE_TYPE_MULTI,
- queue_callback, NULL, UINT32_MAX, UINT32_MAX,
- &agent->kernel_dispatch_command_q);
+ 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_agent_iterate_regions (agent->id, get_kernarg_memory_region,
- &agent->kernarg_region);
+ 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 "
HSA_DEBUG ("Destroying the current HSA program.\n");
- status = hsa_executable_destroy (agent->executable);
+ status = hsa_fns.hsa_executable_destroy_fn (agent->executable);
if (status != HSA_STATUS_SUCCESS)
return hsa_error ("Could not destroy HSA executable", status);
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. */
pair->end = (uintptr_t) (kernel + 1);
struct hsa_kernel_description *d = &image_desc->kernel_infos[i];
- 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 -1;
- }
-
+ if (!init_basic_kernel_info (kernel, d, agent, module))
+ return -1;
kernel++;
pair++;
}
if (agent->prog_finalized)
goto final;
- status = hsa_ext_program_create (HSA_MACHINE_MODEL_LARGE, HSA_PROFILE_FULL,
- HSA_DEFAULT_FLOAT_ROUNDING_MODE_DEFAULT,
- NULL, &prog_handle);
+ 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);
struct module_info *module = agent->first_module;
while (module)
{
- status = hsa_ext_program_add_module (prog_handle,
- module->image_desc->brig_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;
continue;
}
- status = hsa_ext_program_add_module (prog_handle, library->image);
+ 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);
hsa_ext_control_directives_t control_directives;
memset (&control_directives, 0, sizeof (control_directives));
hsa_code_object_t code_object;
- status = hsa_ext_program_finalize (prog_handle, agent->isa,
- HSA_EXT_FINALIZER_CALL_CONVENTION_AUTO,
- control_directives, "",
- HSA_CODE_OBJECT_TYPE_PROGRAM,
- &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);
}
HSA_DEBUG ("Finalization done\n");
- hsa_ext_program_destroy (prog_handle);
+ hsa_fns.hsa_ext_program_destroy_fn (prog_handle);
status
- = hsa_executable_create (HSA_PROFILE_FULL, HSA_EXECUTABLE_STATE_UNFROZEN,
- "", &agent->executable);
+ = 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);
{
struct global_var_info *var;
var = &module->image_desc->global_variables[i];
- status
- = hsa_executable_global_variable_define (agent->executable,
- var->name, var->address);
+ 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);
module = module->next;
}
- status = hsa_executable_load_code_object (agent->executable, agent->id,
- code_object, "");
+ 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_executable_freeze (agent->executable, "");
+ status = hsa_fns.hsa_executable_freeze_fn (agent->executable, "");
if (status != HSA_STATUS_SUCCESS)
hsa_fatal ("Could not freeze the HSA executable", status);
shadow->object = kernel->object;
hsa_signal_t sync_signal;
- hsa_status_t status = hsa_signal_create (1, 0, NULL, &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->group_segment_size = kernel->group_segment_size;
status
- = hsa_memory_allocate (agent->kernarg_region, kernel->kernarg_segment_size,
- &shadow->kernarg_address);
+ = 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);
HSA_DEBUG ("Released kernel dispatch: %p has value: %lu (%p)\n", shadow,
shadow->debug, (void *) shadow->debug);
- hsa_memory_free (shadow->kernarg_address);
+ hsa_fns.hsa_memory_free_fn (shadow->kernarg_address);
hsa_signal_t s;
s.handle = shadow->signal;
- hsa_signal_destroy (s);
+ hsa_fns.hsa_signal_destroy_fn (s);
free (shadow->omp_data_memory);
hsa_status_t status;
struct agent_info *agent = kernel->agent;
hsa_executable_symbol_t kernel_symbol;
- status = hsa_executable_get_symbol (agent->executable, NULL, kernel->name,
- agent->id, 0, &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_executable_symbol_get_info (kernel_symbol,
- HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT,
- &kernel->object);
+ 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_executable_symbol_get_info
+ 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_executable_symbol_get_info
+ 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_executable_symbol_get_info
+ 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)
struct GOMP_kernel_launch_attributes *kla;
kla = (struct GOMP_kernel_launch_attributes *) *input;
*result = kla;
- if (kla->ndim != 1)
- GOMP_PLUGIN_fatal ("HSA does not yet support number of dimensions "
- "different from one.");
- if (kla->gdims[0] == 0)
- return false;
-
- HSA_DEBUG ("GOMP_OFFLOAD_run called with grid size %u and group size %u\n",
- kla->gdims[0], kla->wdims[0]);
+ 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
__atomic_store_n (packet, header | (rest << 16), __ATOMIC_RELEASE);
}
-/* Part of the libgomp plugin interface. Run a kernel on device N 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. */
+/* Run KERNEL on its agent, pass VARS to it as arguments and take
+ launchattributes from KLA. */
void
-GOMP_OFFLOAD_run (int n, void *fn_ptr, void *vars, void **args)
+run_kernel (struct kernel_info *kernel, void *vars,
+ struct GOMP_kernel_launch_attributes *kla)
{
- struct kernel_info *kernel = (struct kernel_info *) fn_ptr;
struct agent_info *agent = kernel->agent;
- 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;
- }
if (pthread_rwlock_rdlock (&agent->modules_rwlock))
GOMP_PLUGIN_fatal ("Unable to read-lock an HSA agent rwlock");
print_kernel_dispatch (shadow, 2);
}
- uint64_t index = hsa_queue_add_write_index_release (agent->command_q, 1);
+ 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_queue_load_read_index_acquire (agent->command_q)
+ while (index - hsa_fns.hsa_queue_load_read_index_acquire_fn (agent->command_q)
>= agent->command_q->size)
;
memset (((uint8_t *) packet) + 4, 0, sizeof (*packet) - 4);
packet->grid_size_x = kla->gdims[0];
- uint32_t wgs = kla->wdims[0];
- if (wgs == 0)
- /* TODO: Provide a default via environment. */
- wgs = 64;
- else if (wgs > kla->gdims[0])
- wgs = kla->gdims[0];
- packet->workgroup_size_x = wgs;
- packet->grid_size_y = 1;
- packet->workgroup_size_y = 1;
- packet->grid_size_z = 1;
- packet->workgroup_size_z = 1;
+ 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;
hsa_signal_t s;
s.handle = shadow->signal;
packet->completion_signal = s;
- hsa_signal_store_relaxed (s, 1);
+ hsa_fns.hsa_signal_store_relaxed_fn (s, 1);
memcpy (shadow->kernarg_address, &vars, sizeof (vars));
/* PR hsa/70337. */
HSA_DEBUG ("Going to dispatch kernel %s\n", kernel->name);
packet_store_release ((uint32_t *) packet, header,
- 1 << HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS);
+ (uint16_t) kla->ndim << HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS);
- hsa_signal_store_release (agent->command_q->doorbell_signal, index);
+ 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
HSA_DEBUG ("Kernel dispatched, waiting for completion\n");
/* Root signal waits with 1ms timeout. */
- while (hsa_signal_wait_acquire (s, HSA_SIGNAL_CONDITION_LT, 1, 1000 * 1000,
- HSA_WAIT_STATE_BLOCKED) != 0)
+ 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;
HSA_DEBUG ("Waiting for children completion signal: %lu\n",
shadow->children_dispatches[i]->signal);
- hsa_signal_load_acquire (child_s);
+ hsa_fns.hsa_signal_load_acquire_fn (child_s);
}
release_kernel_dispatch (shadow);
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
release_agent_shared_libraries (agent);
- hsa_status_t status = hsa_queue_destroy (agent->command_q);
+ 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_queue_destroy (agent->kernel_dispatch_command_q);
+ 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))
-# Makefile.in generated by automake 1.11.6 from Makefile.am.
+# Makefile.in generated by automake 1.11.1 from Makefile.am.
# @configure_input@
# Copyright (C) 1994, 1995, 1996, 1997, 1998, 1999, 2000, 2001, 2002,
-# 2003, 2004, 2005, 2006, 2007, 2008, 2009, 2010, 2011 Free Software
-# Foundation, Inc.
+# 2003, 2004, 2005, 2006, 2007, 2008, 2009 Free Software Foundation,
+# Inc.
# This Makefile.in is free software; the Free Software Foundation
# gives unlimited permission to copy and/or distribute it,
# with or without modifications, as long as this notice is preserved.
@SET_MAKE@
VPATH = @srcdir@
-am__make_dryrun = \
- { \
- am__dry=no; \
- case $$MAKEFLAGS in \
- *\\[\ \ ]*) \
- echo 'am--echo: ; @echo "AM" OK' | $(MAKE) -f - 2>/dev/null \
- | grep '^AM OK$$' >/dev/null || am__dry=yes;; \
- *) \
- for am__flg in $$MAKEFLAGS; do \
- case $$am__flg in \
- *=*|--*) ;; \
- *n*) am__dry=yes; break;; \
- esac; \
- done;; \
- esac; \
- test $$am__dry = yes; \
- }
pkgdatadir = $(datadir)/@PACKAGE@
pkgincludedir = $(includedir)/@PACKAGE@
pkglibdir = $(libdir)/@PACKAGE@
CONFIG_CLEAN_FILES = libgomp-test-support.pt.exp
CONFIG_CLEAN_VPATH_FILES =
SOURCES =
-am__can_run_installinfo = \
- case $$AM_UPDATE_INFO_DIR in \
- n|no|NO) false;; \
- *) (install-info --version) >/dev/null 2>&1;; \
- esac
DEJATOOL = $(PACKAGE)
RUNTESTDEFAULTFLAGS = --tool $$tool --srcdir $$srcdir
ACLOCAL = @ACLOCAL@
FCFLAGS = @FCFLAGS@
FGREP = @FGREP@
GREP = @GREP@
-HSA_KMT_LIB = @HSA_KMT_LIB@
HSA_RUNTIME_INCLUDE = @HSA_RUNTIME_INCLUDE@
HSA_RUNTIME_LIB = @HSA_RUNTIME_LIB@
INSTALL = @INSTALL@
check-DEJAGNU: site.exp
- srcdir='$(srcdir)'; export srcdir; \
+ srcdir=`$(am__cd) $(srcdir) && pwd`; export srcdir; \
EXPECT=$(EXPECT); export EXPECT; \
runtest=$(RUNTEST); \
if $(SHELL) -c "$$runtest --version" > /dev/null 2>&1; then \
else echo "WARNING: could not find \`runtest'" 1>&2; :;\
fi; \
exit $$exit_status
-site.exp: Makefile $(EXTRA_DEJAGNU_SITE_CONFIG)
+site.exp: Makefile
@echo 'Making a new site.exp file...'
@echo '## these variables are automatically generated by make ##' >site.tmp
@echo '# Do not edit here. If you wish to override these values' >>site.tmp
@echo '# edit the last section' >>site.tmp
- @echo 'set srcdir "$(srcdir)"' >>site.tmp
+ @echo 'set srcdir $(srcdir)' >>site.tmp
@echo "set objdir `pwd`" >>site.tmp
@echo 'set build_alias "$(build_alias)"' >>site.tmp
@echo 'set build_triplet $(build_triplet)' >>site.tmp
@echo 'set host_triplet $(host_triplet)' >>site.tmp
@echo 'set target_alias "$(target_alias)"' >>site.tmp
@echo 'set target_triplet $(target_triplet)' >>site.tmp
- @list='$(EXTRA_DEJAGNU_SITE_CONFIG)'; for f in $$list; do \
- echo "## Begin content included from file $$f. Do not modify. ##" \
- && cat `test -f "$$f" || echo '$(srcdir)/'`$$f \
- && echo "## End content included from file $$f. ##" \
- || exit 1; \
- done >> site.tmp
- @echo "## End of auto-generated content; you can edit from here. ##" >> site.tmp
- @if test -f site.exp; then \
- sed -e '1,/^## End of auto-generated content.*##/d' site.exp >> site.tmp; \
- fi
+ @echo '## All variables above are generated by configure. Do Not Edit ##' >>site.tmp
+ @test ! -f site.exp || \
+ sed '1,/^## All variables above are.*##/ d' site.exp >> site.tmp
@-rm -f site.bak
@test ! -f site.exp || mv site.exp site.bak
@mv site.tmp site.exp
installcheck: installcheck-am
install-strip:
- if test -z '$(STRIP)'; then \
- $(MAKE) $(AM_MAKEFLAGS) INSTALL_PROGRAM="$(INSTALL_STRIP_PROGRAM)" \
- install_sh_PROGRAM="$(INSTALL_STRIP_PROGRAM)" INSTALL_STRIP_FLAG=-s \
- install; \
- else \
- $(MAKE) $(AM_MAKEFLAGS) INSTALL_PROGRAM="$(INSTALL_STRIP_PROGRAM)" \
- install_sh_PROGRAM="$(INSTALL_STRIP_PROGRAM)" INSTALL_STRIP_FLAG=-s \
- "INSTALL_PROGRAM_ENV=STRIPPROG='$(STRIP)'" install; \
- fi
+ $(MAKE) $(AM_MAKEFLAGS) INSTALL_PROGRAM="$(INSTALL_STRIP_PROGRAM)" \
+ install_sh_PROGRAM="$(INSTALL_STRIP_PROGRAM)" INSTALL_STRIP_FLAG=-s \
+ `test -z '$(STRIP)' || \
+ echo "INSTALL_PROGRAM_ENV=STRIPPROG='$(STRIP)'"` install
mostlyclean-generic:
clean-generic:
append always_ld_library_path ":$cuda_driver_lib"
}
global hsa_runtime_lib
- global hsa_kmt_lib
if { $hsa_runtime_lib != "" } {
append always_ld_library_path ":$hsa_runtime_lib"
}
- if { $hsa_kmt_lib != "" } {
- append always_ld_library_path ":$hsa_kmt_lib"
- }
}
# We use atomic operations in the testcases to validate results.
set cuda_driver_include "@CUDA_DRIVER_INCLUDE@"
set cuda_driver_lib "@CUDA_DRIVER_LIB@"
set hsa_runtime_lib "@HSA_RUNTIME_LIB@"
-set hsa_kmt_lib "@HSA_KMT_LIB@"
set offload_targets "@offload_targets@"
projects / gcc.git / commitdiff