OpenE2K
/
gcc
2
0
Fork 0
Code Issues Pull Requests Projects Releases Wiki Activity

Remove build dependence on HSA run-time 

2016-11-23  Martin Liska  <mliska@suse.cz>
            Martin Jambor  <mjambor@suse.cz>

gcc/
	* doc/install.texi: Remove entry about --with-hsa-kmt-lib.

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



Co-Authored-By: Martin Jambor <mjambor@suse.cz>

From-SVN: r242749
This commit is contained in:
Martin Liska 2016-11-23 13:27:13 +01:00 committed by Martin Jambor
parent 3615816da8
commit b8d89b03db
15 changed files with 1539 additions and 388 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

5
gcc/ChangeLog
View File

@ -1,3 +1,8 @@
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

6
gcc/doc/install.texi
View File

@ -2035,12 +2035,6 @@ explicitly specify the directory where they are installed. The
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

50
libgomp/ChangeLog
View File

@ -1,3 +1,53 @@
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>

138
libgomp/Makefile.in
View File

@ -1,9 +1,9 @@
# 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.
@ -45,23 +45,6 @@
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@
@ -137,12 +120,6 @@ am__nobase_list = $(am__nobase_strip_setup); \
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)"
@ -226,11 +203,6 @@ RECURSIVE_TARGETS = all-recursive check-recursive dvi-recursive \
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 \
@ -268,7 +240,6 @@ FC = @FC@
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@
@ -450,7 +421,9 @@ libgomp_la_SOURCES = alloc.c barrier.c critical.c env.c error.c iter.c \
# 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)
@ -491,7 +464,7 @@ all: config.h
.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 \
@ -516,7 +489,6 @@ Makefile: $(srcdir)/Makefile.in $(top_builddir)/config.status
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
@ -528,8 +500,10 @@ $(ACLOCAL_M4): @MAINTAINER_MODE_TRUE@ $(am__aclocal_m4_deps)
$(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
@ -553,6 +527,7 @@ libgomp.spec: $(top_builddir)/config.status $(srcdir)/libgomp.spec.in
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 \
@ -560,8 +535,6 @@ install-toolexeclibLTLIBRARIES: $(toolexeclib_LTLIBRARIES)
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)"; \
}
@ -583,11 +556,11 @@ clean-toolexeclibLTLIBRARIES:
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:
@ -752,7 +725,9 @@ uninstall-html-am:
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|^.*/||'`; \
@ -826,11 +801,8 @@ maintainer-clean-aminfo:
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"; \
@ -844,14 +816,13 @@ uninstall-nodist_fincludeHEADERS:
@$(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"; \
@ -865,14 +836,13 @@ uninstall-nodist_libsubincludeHEADERS:
@$(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"; \
@ -886,7 +856,9 @@ uninstall-nodist_toolexeclibHEADERS:
@$(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.
@ -1041,15 +1013,10 @@ install-am: all-am
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:
@ -1094,11 +1061,8 @@ install-dvi: install-dvi-recursive
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"; \
@ -1114,22 +1078,18 @@ install-html: install-html-recursive
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) | \
@ -1141,12 +1101,9 @@ install-info: install-info-recursive
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/||"`;; \
@ -1164,7 +1121,8 @@ install-info-am: $(INFO_DEPS)
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|^.*/||'`; \
@ -1178,11 +1136,8 @@ install-pdf: install-pdf-recursive
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"; \
@ -1194,11 +1149,8 @@ install-ps: install-ps-recursive
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"; \

74
libgomp/aclocal.m4 vendored
View File

@ -1,8 +1,7 @@
# 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.
@ -20,15 +19,12 @@ You have another version of autoconf. It may work, but is not guaranteed to.
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
@ -38,7 +34,7 @@ AC_DEFUN([AM_AUTOMAKE_VERSION],
[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
])
@ -54,21 +50,19 @@ m4_define([_AM_AUTOCONF_VERSION], [])
# 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/../..'.
@ -150,14 +144,14 @@ AC_CONFIG_COMMANDS_PRE(
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,
@ -197,7 +191,6 @@ AC_CACHE_CHECK([dependency style of $depcc],
# 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.
@ -262,7 +255,7 @@ AC_CACHE_CHECK([dependency style of $depcc],
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.
@ -327,13 +320,10 @@ AC_DEFUN([AM_DEP_TRACK],
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 -*-
@ -555,15 +545,12 @@ for _am_header in $config_headers :; do
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.
@ -582,8 +569,8 @@ AC_SUBST(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,
@ -603,7 +590,7 @@ AC_DEFUN([AM_MAINTAINER_MODE],
[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
@ -714,15 +701,12 @@ else
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'.
@ -745,14 +729,13 @@ esac
# 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)
# -----------------------
@ -760,13 +743,13 @@ AC_DEFUN([_AM_MANGLE_OPTION],
[[_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)])])
@ -842,14 +825,12 @@ Check your system clock])
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
@ -872,13 +853,13 @@ fi
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)
# ---------------------------
@ -887,13 +868,13 @@ AC_SUBST([INSTALL_STRIP_PROGRAM])])
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,
@ -915,11 +896,10 @@ AC_DEFUN([AM_SUBST_NOTMAKE], [_AM_SUBST_NOTMAKE($@)])
# 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])

21
libgomp/config.h.in
View File

@ -24,9 +24,21 @@
/* 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
@ -42,6 +54,9 @@
/* 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
@ -91,6 +106,12 @@
/* 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

129
libgomp/configure vendored
View File

@ -597,6 +597,8 @@ ac_includes_default="\
# include <unistd.h>
#endif"
ac_header_list=
ac_func_list=
ac_subst_vars='am__EXEEXT_FALSE
am__EXEEXT_TRUE
LTLIBOBJS
@ -637,7 +639,6 @@ PLUGIN_HSA_LIBS
PLUGIN_HSA_LDFLAGS
PLUGIN_HSA_CPPFLAGS
PLUGIN_HSA
HSA_KMT_LIB
HSA_RUNTIME_LIB
HSA_RUNTIME_INCLUDE
PLUGIN_NVPTX_LIBS
@ -682,7 +683,6 @@ AR
am__fastdepCC_FALSE
am__fastdepCC_TRUE
CCDEPMODE
am__nodep
AMDEPBACKSLASH
AMDEP_FALSE
AMDEP_TRUE
@ -794,7 +794,6 @@ with_cuda_driver_lib
with_hsa_runtime
with_hsa_runtime_include
with_hsa_runtime_lib
with_hsa_kmt_lib
enable_linux_futex
enable_tls
enable_symvers
@ -1476,7 +1475,6 @@ Optional Packages:
--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
@ -2518,6 +2516,13 @@ $as_echo "$as_me: creating cache $cache_file" >&6;}
>$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
@ -3280,11 +3285,11 @@ MAKEINFO=${MAKEINFO-"${am_missing_run}makeinfo"}
# 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 -'
@ -4182,7 +4187,6 @@ fi
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=
@ -4207,7 +4211,6 @@ else
# 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.
@ -4267,7 +4270,7 @@ else
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.
@ -11145,7 +11148,7 @@ else
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
@ -11251,7 +11254,7 @@ else
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
@ -15198,6 +15201,57 @@ elif test "x${enable_offload_targets-no}" != xno; then
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=
@ -15293,22 +15347,6 @@ if test "x$HSA_RUNTIME_LIB" != x; then
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=
@ -15318,8 +15356,6 @@ PLUGIN_HSA_LIBS=
# Get offload targets and path to install tree of offloading compiler.
offload_additional_options=
offload_additional_lib_paths=
@ -15384,8 +15420,8 @@ rm -f core conftest.err conftest.$ac_objext \
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"
@ -15394,22 +15430,7 @@ rm -f core conftest.err conftest.$ac_objext \
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
@ -15484,6 +15505,16 @@ cat >>confdefs.h <<_ACEOF
_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

3
libgomp/plugin/Makefrag.am
View File

@ -44,7 +44,8 @@ if PLUGIN_HSA
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)

35
libgomp/plugin/configfrag.ac
View File

@ -36,6 +36,9 @@ if test x"$plugin_support" = xyes; then
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=
@ -118,19 +121,6 @@ if test "x$HSA_RUNTIME_LIB" != x; then
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=
@ -140,8 +130,6 @@ AC_SUBST(PLUGIN_HSA_CPPFLAGS)
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=
@ -195,8 +183,8 @@ if test x"$enable_offload_targets" != x; then
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"
@ -205,11 +193,7 @@ if test x"$enable_offload_targets" != x; then
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
@ -260,3 +244,10 @@ AC_DEFINE_UNQUOTED([PLUGIN_NVPTX], [$PLUGIN_NVPTX],
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.])

630
libgomp/plugin/hsa.h Normal file
View File

@ -0,0 +1,630 @@
/* 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 */

265
libgomp/plugin/hsa_ext_finalize.h Normal file
View File

@ -0,0 +1,265 @@
/* 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 */

505
libgomp/plugin/plugin-hsa.c
View File

@ -27,16 +27,129 @@
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. */
@ -129,20 +242,36 @@ static bool debug;
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
@ -176,7 +305,7 @@ hsa_warn (const char *str, hsa_status_t status)
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);
}
@ -188,7 +317,7 @@ static void
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);
}
@ -200,7 +329,7 @@ static bool
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;
@ -359,6 +488,50 @@ struct hsa_context_info
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 *
@ -386,17 +559,32 @@ suitable_hsa_agent_p (hsa_agent_t agent)
{
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;
@ -443,11 +631,16 @@ init_hsa_context (void)
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);
@ -455,7 +648,7 @@ init_hsa_context (void)
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");
@ -485,14 +678,16 @@ get_kernarg_memory_region (hsa_region_t region, void *data)
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)
@ -546,29 +741,36 @@ GOMP_OFFLOAD_init_device (int n)
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 "
@ -646,7 +848,7 @@ destroy_hsa_program (struct agent_info *agent)
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);
@ -661,6 +863,29 @@ destroy_hsa_program (struct agent_info *agent)
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. */
@ -715,19 +940,8 @@ GOMP_OFFLOAD_load_image (int ord, unsigned version, void *target_data,
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++;
}
@ -799,9 +1013,10 @@ create_and_finalize_hsa_program (struct agent_info *agent)
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);
@ -810,8 +1025,8 @@ create_and_finalize_hsa_program (struct agent_info *agent)
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;
@ -837,7 +1052,8 @@ create_and_finalize_hsa_program (struct agent_info *agent)
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);
@ -849,11 +1065,9 @@ create_and_finalize_hsa_program (struct agent_info *agent)
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);
@ -861,11 +1075,12 @@ create_and_finalize_hsa_program (struct agent_info *agent)
}
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);
@ -877,9 +1092,8 @@ create_and_finalize_hsa_program (struct agent_info *agent)
{
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);
@ -892,11 +1106,12 @@ create_and_finalize_hsa_program (struct agent_info *agent)
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);
@ -937,7 +1152,7 @@ create_single_kernel_dispatch (struct kernel_info *kernel,
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);
@ -946,8 +1161,9 @@ create_single_kernel_dispatch (struct kernel_info *kernel,
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);
@ -962,11 +1178,11 @@ release_kernel_dispatch (struct GOMP_hsa_kernel_dispatch *shadow)
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);
@ -986,31 +1202,30 @@ init_single_kernel (struct kernel_info *kernel, unsigned *max_omp_data_size)
hsa_status_t status;
struct agent_info *agent = kernel->agent;
hsa_executable_symbol_t kernel_symbol;
status = hsa_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)
@ -1209,18 +1424,43 @@ parse_target_attributes (void **input,
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
@ -1254,22 +1494,14 @@ packet_store_release (uint32_t* packet, uint16_t header, uint16_t rest)
__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");
@ -1288,11 +1520,12 @@ GOMP_OFFLOAD_run (int n, void *fn_ptr, void *vars, void **args)
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)
;
@ -1302,17 +1535,33 @@ GOMP_OFFLOAD_run (int n, void *fn_ptr, void *vars, void **args)
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;
@ -1320,7 +1569,7 @@ GOMP_OFFLOAD_run (int n, void *fn_ptr, void *vars, void **args)
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. */
@ -1344,9 +1593,10 @@ GOMP_OFFLOAD_run (int n, void *fn_ptr, void *vars, void **args)
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
@ -1357,8 +1607,9 @@ GOMP_OFFLOAD_run (int n, void *fn_ptr, void *vars, void **args)
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;
@ -1366,7 +1617,7 @@ GOMP_OFFLOAD_run (int n, void *fn_ptr, void *vars, void **args)
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);
@ -1375,6 +1626,26 @@ GOMP_OFFLOAD_run (int n, void *fn_ptr, void *vars, void **args)
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
@ -1534,10 +1805,10 @@ GOMP_OFFLOAD_fini_device (int n)
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))

61
libgomp/testsuite/Makefile.in
View File

@ -1,9 +1,9 @@
# 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.
@ -15,23 +15,6 @@
@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@
@ -76,11 +59,6 @@ CONFIG_HEADER = $(top_builddir)/config.h
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@
@ -111,7 +89,6 @@ FC = @FC@
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@
@ -303,7 +280,7 @@ CTAGS:
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 \
@ -314,12 +291,12 @@ check-DEJAGNU: site.exp
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
@ -327,16 +304,9 @@ site.exp: Makefile $(EXTRA_DEJAGNU_SITE_CONFIG)
@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
@ -361,15 +331,10 @@ install-am: all-am
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:

4
libgomp/testsuite/lib/libgomp.exp
View File

@ -205,13 +205,9 @@ proc libgomp_init { args } {
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.

1
libgomp/testsuite/libgomp-test-support.exp.in
View File

@ -1,6 +1,5 @@
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@"