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

Remove libhsail-rt. 

ChangeLog:

	* Makefile.def: Remove libhsail-rt.
	* Makefile.in: Likewise.
	* configure.ac: Likewise.
	* configure: Regenerate.
	* libhsail-rt/ChangeLog: Removed.
	* libhsail-rt/Makefile.am: Removed.
	* libhsail-rt/Makefile.in: Removed.
	* libhsail-rt/README: Removed.
	* libhsail-rt/aclocal.m4: Removed.
	* libhsail-rt/configure: Removed.
	* libhsail-rt/configure.ac: Removed.
	* libhsail-rt/configure.tgt: Removed.
	* libhsail-rt/include/internal/fibers.h: Removed.
	* libhsail-rt/include/internal/phsa-queue-interface.h: Removed.
	* libhsail-rt/include/internal/phsa-rt.h: Removed.
	* libhsail-rt/include/internal/workitems.h: Removed.
	* libhsail-rt/rt/arithmetic.c: Removed.
	* libhsail-rt/rt/atomics.c: Removed.
	* libhsail-rt/rt/bitstring.c: Removed.
	* libhsail-rt/rt/fbarrier.c: Removed.
	* libhsail-rt/rt/fibers.c: Removed.
	* libhsail-rt/rt/fp16.c: Removed.
	* libhsail-rt/rt/misc.c: Removed.
	* libhsail-rt/rt/multimedia.c: Removed.
	* libhsail-rt/rt/queue.c: Removed.
	* libhsail-rt/rt/sat_arithmetic.c: Removed.
	* libhsail-rt/rt/segment.c: Removed.
	* libhsail-rt/rt/workitems.c: Removed.
	* libhsail-rt/target-config.h.in: Removed.

contrib/ChangeLog:

	* gcc_update: Remove libhsail-rt folder.
	* update-copyright.py: Likewise.
This commit is contained in:
Martin Liska 2021-03-11 16:23:56 +01:00
parent 900b1c27b9
commit 814d86ddce
31 changed files with 0 additions and 23641 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

3
Makefile.def
View File

@ -166,7 +166,6 @@ target_modules = { module= libquadmath; };
target_modules = { module= libgfortran; };
target_modules = { module= libobjc; };
target_modules = { module= libgo; };
target_modules = { module= libhsail-rt; };
target_modules = { module= libphobos;
lib_path=src/.libs; };
target_modules = { module= libtermcap; no_check=true;
@ -651,8 +650,6 @@ languages = { language=obj-c++; gcc-check-target=check-obj-c++; };
languages = { language=go; gcc-check-target=check-go;
lib-check-target=check-target-libgo;
lib-check-target=check-gotools; };
languages = { language=brig; gcc-check-target=check-brig;
lib-check-target=check-target-libhsail-rt; };
languages = { language=d; gcc-check-target=check-d;
lib-check-target=check-target-libphobos; };

489
Makefile.in
View File

@ -1058,7 +1058,6 @@ configure-target: \
maybe-configure-target-libgfortran \
maybe-configure-target-libobjc \
maybe-configure-target-libgo \
maybe-configure-target-libhsail-rt \
maybe-configure-target-libphobos \
maybe-configure-target-libtermcap \
maybe-configure-target-winsup \
@ -1231,7 +1230,6 @@ all-target: maybe-all-target-libquadmath
all-target: maybe-all-target-libgfortran
all-target: maybe-all-target-libobjc
all-target: maybe-all-target-libgo
all-target: maybe-all-target-libhsail-rt
all-target: maybe-all-target-libphobos
all-target: maybe-all-target-libtermcap
all-target: maybe-all-target-winsup
@ -1329,7 +1327,6 @@ info-target: maybe-info-target-libquadmath
info-target: maybe-info-target-libgfortran
info-target: maybe-info-target-libobjc
info-target: maybe-info-target-libgo
info-target: maybe-info-target-libhsail-rt
info-target: maybe-info-target-libphobos
info-target: maybe-info-target-libtermcap
info-target: maybe-info-target-winsup
@ -1420,7 +1417,6 @@ dvi-target: maybe-dvi-target-libquadmath
dvi-target: maybe-dvi-target-libgfortran
dvi-target: maybe-dvi-target-libobjc
dvi-target: maybe-dvi-target-libgo
dvi-target: maybe-dvi-target-libhsail-rt
dvi-target: maybe-dvi-target-libphobos
dvi-target: maybe-dvi-target-libtermcap
dvi-target: maybe-dvi-target-winsup
@ -1511,7 +1507,6 @@ pdf-target: maybe-pdf-target-libquadmath
pdf-target: maybe-pdf-target-libgfortran
pdf-target: maybe-pdf-target-libobjc
pdf-target: maybe-pdf-target-libgo
pdf-target: maybe-pdf-target-libhsail-rt
pdf-target: maybe-pdf-target-libphobos
pdf-target: maybe-pdf-target-libtermcap
pdf-target: maybe-pdf-target-winsup
@ -1602,7 +1597,6 @@ html-target: maybe-html-target-libquadmath
html-target: maybe-html-target-libgfortran
html-target: maybe-html-target-libobjc
html-target: maybe-html-target-libgo
html-target: maybe-html-target-libhsail-rt
html-target: maybe-html-target-libphobos
html-target: maybe-html-target-libtermcap
html-target: maybe-html-target-winsup
@ -1693,7 +1687,6 @@ TAGS-target: maybe-TAGS-target-libquadmath
TAGS-target: maybe-TAGS-target-libgfortran
TAGS-target: maybe-TAGS-target-libobjc
TAGS-target: maybe-TAGS-target-libgo
TAGS-target: maybe-TAGS-target-libhsail-rt
TAGS-target: maybe-TAGS-target-libphobos
TAGS-target: maybe-TAGS-target-libtermcap
TAGS-target: maybe-TAGS-target-winsup
@ -1784,7 +1777,6 @@ install-info-target: maybe-install-info-target-libquadmath
install-info-target: maybe-install-info-target-libgfortran
install-info-target: maybe-install-info-target-libobjc
install-info-target: maybe-install-info-target-libgo
install-info-target: maybe-install-info-target-libhsail-rt
install-info-target: maybe-install-info-target-libphobos
install-info-target: maybe-install-info-target-libtermcap
install-info-target: maybe-install-info-target-winsup
@ -1875,7 +1867,6 @@ install-pdf-target: maybe-install-pdf-target-libquadmath
install-pdf-target: maybe-install-pdf-target-libgfortran
install-pdf-target: maybe-install-pdf-target-libobjc
install-pdf-target: maybe-install-pdf-target-libgo
install-pdf-target: maybe-install-pdf-target-libhsail-rt
install-pdf-target: maybe-install-pdf-target-libphobos
install-pdf-target: maybe-install-pdf-target-libtermcap
install-pdf-target: maybe-install-pdf-target-winsup
@ -1966,7 +1957,6 @@ install-html-target: maybe-install-html-target-libquadmath
install-html-target: maybe-install-html-target-libgfortran
install-html-target: maybe-install-html-target-libobjc
install-html-target: maybe-install-html-target-libgo
install-html-target: maybe-install-html-target-libhsail-rt
install-html-target: maybe-install-html-target-libphobos
install-html-target: maybe-install-html-target-libtermcap
install-html-target: maybe-install-html-target-winsup
@ -2057,7 +2047,6 @@ installcheck-target: maybe-installcheck-target-libquadmath
installcheck-target: maybe-installcheck-target-libgfortran
installcheck-target: maybe-installcheck-target-libobjc
installcheck-target: maybe-installcheck-target-libgo
installcheck-target: maybe-installcheck-target-libhsail-rt
installcheck-target: maybe-installcheck-target-libphobos
installcheck-target: maybe-installcheck-target-libtermcap
installcheck-target: maybe-installcheck-target-winsup
@ -2148,7 +2137,6 @@ mostlyclean-target: maybe-mostlyclean-target-libquadmath
mostlyclean-target: maybe-mostlyclean-target-libgfortran
mostlyclean-target: maybe-mostlyclean-target-libobjc
mostlyclean-target: maybe-mostlyclean-target-libgo
mostlyclean-target: maybe-mostlyclean-target-libhsail-rt
mostlyclean-target: maybe-mostlyclean-target-libphobos
mostlyclean-target: maybe-mostlyclean-target-libtermcap
mostlyclean-target: maybe-mostlyclean-target-winsup
@ -2239,7 +2227,6 @@ clean-target: maybe-clean-target-libquadmath
clean-target: maybe-clean-target-libgfortran
clean-target: maybe-clean-target-libobjc
clean-target: maybe-clean-target-libgo
clean-target: maybe-clean-target-libhsail-rt
clean-target: maybe-clean-target-libphobos
clean-target: maybe-clean-target-libtermcap
clean-target: maybe-clean-target-winsup
@ -2330,7 +2317,6 @@ distclean-target: maybe-distclean-target-libquadmath
distclean-target: maybe-distclean-target-libgfortran
distclean-target: maybe-distclean-target-libobjc
distclean-target: maybe-distclean-target-libgo
distclean-target: maybe-distclean-target-libhsail-rt
distclean-target: maybe-distclean-target-libphobos
distclean-target: maybe-distclean-target-libtermcap
distclean-target: maybe-distclean-target-winsup
@ -2421,7 +2407,6 @@ maintainer-clean-target: maybe-maintainer-clean-target-libquadmath
maintainer-clean-target: maybe-maintainer-clean-target-libgfortran
maintainer-clean-target: maybe-maintainer-clean-target-libobjc
maintainer-clean-target: maybe-maintainer-clean-target-libgo
maintainer-clean-target: maybe-maintainer-clean-target-libhsail-rt
maintainer-clean-target: maybe-maintainer-clean-target-libphobos
maintainer-clean-target: maybe-maintainer-clean-target-libtermcap
maintainer-clean-target: maybe-maintainer-clean-target-winsup
@ -2568,7 +2553,6 @@ check-target: \
maybe-check-target-libgfortran \
maybe-check-target-libobjc \
maybe-check-target-libgo \
maybe-check-target-libhsail-rt \
maybe-check-target-libphobos \
maybe-check-target-libtermcap \
maybe-check-target-winsup \
@ -2761,7 +2745,6 @@ install-target: \
maybe-install-target-libgfortran \
maybe-install-target-libobjc \
maybe-install-target-libgo \
maybe-install-target-libhsail-rt \
maybe-install-target-libphobos \
maybe-install-target-libtermcap \
maybe-install-target-winsup \
@ -2872,7 +2855,6 @@ install-strip-target: \
maybe-install-strip-target-libgfortran \
maybe-install-strip-target-libobjc \
maybe-install-strip-target-libgo \
maybe-install-strip-target-libhsail-rt \
maybe-install-strip-target-libphobos \
maybe-install-strip-target-libtermcap \
maybe-install-strip-target-winsup \
@ -50752,464 +50734,6 @@ maintainer-clean-target-libgo:
.PHONY: configure-target-libhsail-rt maybe-configure-target-libhsail-rt
maybe-configure-target-libhsail-rt:
@if gcc-bootstrap
configure-target-libhsail-rt: stage_current
@endif gcc-bootstrap
@if target-libhsail-rt
maybe-configure-target-libhsail-rt: configure-target-libhsail-rt
configure-target-libhsail-rt:
@: $(MAKE); $(unstage)
@r=`${PWD_COMMAND}`; export r; \
s=`cd $(srcdir); ${PWD_COMMAND}`; export s; \
echo "Checking multilib configuration for libhsail-rt..."; \
$(SHELL) $(srcdir)/mkinstalldirs $(TARGET_SUBDIR)/libhsail-rt; \
$(CC_FOR_TARGET) --print-multi-lib > $(TARGET_SUBDIR)/libhsail-rt/multilib.tmp 2> /dev/null; \
if test -r $(TARGET_SUBDIR)/libhsail-rt/multilib.out; then \
if cmp -s $(TARGET_SUBDIR)/libhsail-rt/multilib.tmp $(TARGET_SUBDIR)/libhsail-rt/multilib.out; then \
rm -f $(TARGET_SUBDIR)/libhsail-rt/multilib.tmp; \
else \
rm -f $(TARGET_SUBDIR)/libhsail-rt/Makefile; \
mv $(TARGET_SUBDIR)/libhsail-rt/multilib.tmp $(TARGET_SUBDIR)/libhsail-rt/multilib.out; \
fi; \
else \
mv $(TARGET_SUBDIR)/libhsail-rt/multilib.tmp $(TARGET_SUBDIR)/libhsail-rt/multilib.out; \
fi; \
test ! -f $(TARGET_SUBDIR)/libhsail-rt/Makefile || exit 0; \
$(SHELL) $(srcdir)/mkinstalldirs $(TARGET_SUBDIR)/libhsail-rt; \
$(NORMAL_TARGET_EXPORTS) \
echo Configuring in $(TARGET_SUBDIR)/libhsail-rt; \
cd "$(TARGET_SUBDIR)/libhsail-rt" || exit 1; \
case $(srcdir) in \
/* | [A-Za-z]:[\\/]*) topdir=$(srcdir) ;; \
*) topdir=`echo $(TARGET_SUBDIR)/libhsail-rt/ | \
sed -e 's,\./,,g' -e 's,[^/]*/,../,g' `$(srcdir) ;; \
esac; \
module_srcdir=libhsail-rt; \
rm -f no-such-file || : ; \
CONFIG_SITE=no-such-file $(SHELL) \
$$s/$$module_srcdir/configure \
--srcdir=$${topdir}/$$module_srcdir \
$(TARGET_CONFIGARGS) --build=${build_alias} --host=${target_alias} \
--target=${target_alias} \
|| exit 1
@endif target-libhsail-rt
.PHONY: all-target-libhsail-rt maybe-all-target-libhsail-rt
maybe-all-target-libhsail-rt:
@if gcc-bootstrap
all-target-libhsail-rt: stage_current
@endif gcc-bootstrap
@if target-libhsail-rt
TARGET-target-libhsail-rt=all
maybe-all-target-libhsail-rt: all-target-libhsail-rt
all-target-libhsail-rt: configure-target-libhsail-rt
@: $(MAKE); $(unstage)
@r=`${PWD_COMMAND}`; export r; \
s=`cd $(srcdir); ${PWD_COMMAND}`; export s; \
$(NORMAL_TARGET_EXPORTS) \
(cd $(TARGET_SUBDIR)/libhsail-rt && \
$(MAKE) $(BASE_FLAGS_TO_PASS) $(EXTRA_TARGET_FLAGS) \
$(TARGET-target-libhsail-rt))
@endif target-libhsail-rt
.PHONY: check-target-libhsail-rt maybe-check-target-libhsail-rt
maybe-check-target-libhsail-rt:
@if target-libhsail-rt
maybe-check-target-libhsail-rt: check-target-libhsail-rt
check-target-libhsail-rt:
@: $(MAKE); $(unstage)
@r=`${PWD_COMMAND}`; export r; \
s=`cd $(srcdir); ${PWD_COMMAND}`; export s; \
$(NORMAL_TARGET_EXPORTS) \
(cd $(TARGET_SUBDIR)/libhsail-rt && \
$(MAKE) $(TARGET_FLAGS_TO_PASS) check)
@endif target-libhsail-rt
.PHONY: install-target-libhsail-rt maybe-install-target-libhsail-rt
maybe-install-target-libhsail-rt:
@if target-libhsail-rt
maybe-install-target-libhsail-rt: install-target-libhsail-rt
install-target-libhsail-rt: installdirs
@: $(MAKE); $(unstage)
@r=`${PWD_COMMAND}`; export r; \
s=`cd $(srcdir); ${PWD_COMMAND}`; export s; \
$(NORMAL_TARGET_EXPORTS) \
(cd $(TARGET_SUBDIR)/libhsail-rt && \
$(MAKE) $(TARGET_FLAGS_TO_PASS) install)
@endif target-libhsail-rt
.PHONY: install-strip-target-libhsail-rt maybe-install-strip-target-libhsail-rt
maybe-install-strip-target-libhsail-rt:
@if target-libhsail-rt
maybe-install-strip-target-libhsail-rt: install-strip-target-libhsail-rt
install-strip-target-libhsail-rt: installdirs
@: $(MAKE); $(unstage)
@r=`${PWD_COMMAND}`; export r; \
s=`cd $(srcdir); ${PWD_COMMAND}`; export s; \
$(NORMAL_TARGET_EXPORTS) \
(cd $(TARGET_SUBDIR)/libhsail-rt && \
$(MAKE) $(TARGET_FLAGS_TO_PASS) install-strip)
@endif target-libhsail-rt
# Other targets (info, dvi, pdf, etc.)
.PHONY: maybe-info-target-libhsail-rt info-target-libhsail-rt
maybe-info-target-libhsail-rt:
@if target-libhsail-rt
maybe-info-target-libhsail-rt: info-target-libhsail-rt
info-target-libhsail-rt: \
configure-target-libhsail-rt
@: $(MAKE); $(unstage)
@[ -f $(TARGET_SUBDIR)/libhsail-rt/Makefile ] || exit 0; \
r=`${PWD_COMMAND}`; export r; \
s=`cd $(srcdir); ${PWD_COMMAND}`; export s; \
$(NORMAL_TARGET_EXPORTS) \
echo "Doing info in $(TARGET_SUBDIR)/libhsail-rt"; \
for flag in $(EXTRA_TARGET_FLAGS); do \
eval `echo "$$flag" | sed -e "s|^\([^=]*\)=\(.*\)|\1='\2'; export \1|"`; \
done; \
(cd $(TARGET_SUBDIR)/libhsail-rt && \
$(MAKE) $(BASE_FLAGS_TO_PASS) "AR=$${AR}" "AS=$${AS}" \
"CC=$${CC}" "CXX=$${CXX}" "LD=$${LD}" "NM=$${NM}" \
"RANLIB=$${RANLIB}" \
"DLLTOOL=$${DLLTOOL}" "WINDRES=$${WINDRES}" "WINDMC=$${WINDMC}" \
info) \
|| exit 1
@endif target-libhsail-rt
.PHONY: maybe-dvi-target-libhsail-rt dvi-target-libhsail-rt
maybe-dvi-target-libhsail-rt:
@if target-libhsail-rt
maybe-dvi-target-libhsail-rt: dvi-target-libhsail-rt
dvi-target-libhsail-rt: \
configure-target-libhsail-rt
@: $(MAKE); $(unstage)
@[ -f $(TARGET_SUBDIR)/libhsail-rt/Makefile ] || exit 0; \
r=`${PWD_COMMAND}`; export r; \
s=`cd $(srcdir); ${PWD_COMMAND}`; export s; \
$(NORMAL_TARGET_EXPORTS) \
echo "Doing dvi in $(TARGET_SUBDIR)/libhsail-rt"; \
for flag in $(EXTRA_TARGET_FLAGS); do \
eval `echo "$$flag" | sed -e "s|^\([^=]*\)=\(.*\)|\1='\2'; export \1|"`; \
done; \
(cd $(TARGET_SUBDIR)/libhsail-rt && \
$(MAKE) $(BASE_FLAGS_TO_PASS) "AR=$${AR}" "AS=$${AS}" \
"CC=$${CC}" "CXX=$${CXX}" "LD=$${LD}" "NM=$${NM}" \
"RANLIB=$${RANLIB}" \
"DLLTOOL=$${DLLTOOL}" "WINDRES=$${WINDRES}" "WINDMC=$${WINDMC}" \
dvi) \
|| exit 1
@endif target-libhsail-rt
.PHONY: maybe-pdf-target-libhsail-rt pdf-target-libhsail-rt
maybe-pdf-target-libhsail-rt:
@if target-libhsail-rt
maybe-pdf-target-libhsail-rt: pdf-target-libhsail-rt
pdf-target-libhsail-rt: \
configure-target-libhsail-rt
@: $(MAKE); $(unstage)
@[ -f $(TARGET_SUBDIR)/libhsail-rt/Makefile ] || exit 0; \
r=`${PWD_COMMAND}`; export r; \
s=`cd $(srcdir); ${PWD_COMMAND}`; export s; \
$(NORMAL_TARGET_EXPORTS) \
echo "Doing pdf in $(TARGET_SUBDIR)/libhsail-rt"; \
for flag in $(EXTRA_TARGET_FLAGS); do \
eval `echo "$$flag" | sed -e "s|^\([^=]*\)=\(.*\)|\1='\2'; export \1|"`; \
done; \
(cd $(TARGET_SUBDIR)/libhsail-rt && \
$(MAKE) $(BASE_FLAGS_TO_PASS) "AR=$${AR}" "AS=$${AS}" \
"CC=$${CC}" "CXX=$${CXX}" "LD=$${LD}" "NM=$${NM}" \
"RANLIB=$${RANLIB}" \
"DLLTOOL=$${DLLTOOL}" "WINDRES=$${WINDRES}" "WINDMC=$${WINDMC}" \
pdf) \
|| exit 1
@endif target-libhsail-rt
.PHONY: maybe-html-target-libhsail-rt html-target-libhsail-rt
maybe-html-target-libhsail-rt:
@if target-libhsail-rt
maybe-html-target-libhsail-rt: html-target-libhsail-rt
html-target-libhsail-rt: \
configure-target-libhsail-rt
@: $(MAKE); $(unstage)
@[ -f $(TARGET_SUBDIR)/libhsail-rt/Makefile ] || exit 0; \
r=`${PWD_COMMAND}`; export r; \
s=`cd $(srcdir); ${PWD_COMMAND}`; export s; \
$(NORMAL_TARGET_EXPORTS) \
echo "Doing html in $(TARGET_SUBDIR)/libhsail-rt"; \
for flag in $(EXTRA_TARGET_FLAGS); do \
eval `echo "$$flag" | sed -e "s|^\([^=]*\)=\(.*\)|\1='\2'; export \1|"`; \
done; \
(cd $(TARGET_SUBDIR)/libhsail-rt && \
$(MAKE) $(BASE_FLAGS_TO_PASS) "AR=$${AR}" "AS=$${AS}" \
"CC=$${CC}" "CXX=$${CXX}" "LD=$${LD}" "NM=$${NM}" \
"RANLIB=$${RANLIB}" \
"DLLTOOL=$${DLLTOOL}" "WINDRES=$${WINDRES}" "WINDMC=$${WINDMC}" \
html) \
|| exit 1
@endif target-libhsail-rt
.PHONY: maybe-TAGS-target-libhsail-rt TAGS-target-libhsail-rt
maybe-TAGS-target-libhsail-rt:
@if target-libhsail-rt
maybe-TAGS-target-libhsail-rt: TAGS-target-libhsail-rt
TAGS-target-libhsail-rt: \
configure-target-libhsail-rt
@: $(MAKE); $(unstage)
@[ -f $(TARGET_SUBDIR)/libhsail-rt/Makefile ] || exit 0; \
r=`${PWD_COMMAND}`; export r; \
s=`cd $(srcdir); ${PWD_COMMAND}`; export s; \
$(NORMAL_TARGET_EXPORTS) \
echo "Doing TAGS in $(TARGET_SUBDIR)/libhsail-rt"; \
for flag in $(EXTRA_TARGET_FLAGS); do \
eval `echo "$$flag" | sed -e "s|^\([^=]*\)=\(.*\)|\1='\2'; export \1|"`; \
done; \
(cd $(TARGET_SUBDIR)/libhsail-rt && \
$(MAKE) $(BASE_FLAGS_TO_PASS) "AR=$${AR}" "AS=$${AS}" \
"CC=$${CC}" "CXX=$${CXX}" "LD=$${LD}" "NM=$${NM}" \
"RANLIB=$${RANLIB}" \
"DLLTOOL=$${DLLTOOL}" "WINDRES=$${WINDRES}" "WINDMC=$${WINDMC}" \
TAGS) \
|| exit 1
@endif target-libhsail-rt
.PHONY: maybe-install-info-target-libhsail-rt install-info-target-libhsail-rt
maybe-install-info-target-libhsail-rt:
@if target-libhsail-rt
maybe-install-info-target-libhsail-rt: install-info-target-libhsail-rt
install-info-target-libhsail-rt: \
configure-target-libhsail-rt \
info-target-libhsail-rt
@: $(MAKE); $(unstage)
@[ -f $(TARGET_SUBDIR)/libhsail-rt/Makefile ] || exit 0; \
r=`${PWD_COMMAND}`; export r; \
s=`cd $(srcdir); ${PWD_COMMAND}`; export s; \
$(NORMAL_TARGET_EXPORTS) \
echo "Doing install-info in $(TARGET_SUBDIR)/libhsail-rt"; \
for flag in $(EXTRA_TARGET_FLAGS); do \
eval `echo "$$flag" | sed -e "s|^\([^=]*\)=\(.*\)|\1='\2'; export \1|"`; \
done; \
(cd $(TARGET_SUBDIR)/libhsail-rt && \
$(MAKE) $(BASE_FLAGS_TO_PASS) "AR=$${AR}" "AS=$${AS}" \
"CC=$${CC}" "CXX=$${CXX}" "LD=$${LD}" "NM=$${NM}" \
"RANLIB=$${RANLIB}" \
"DLLTOOL=$${DLLTOOL}" "WINDRES=$${WINDRES}" "WINDMC=$${WINDMC}" \
install-info) \
|| exit 1
@endif target-libhsail-rt
.PHONY: maybe-install-pdf-target-libhsail-rt install-pdf-target-libhsail-rt
maybe-install-pdf-target-libhsail-rt:
@if target-libhsail-rt
maybe-install-pdf-target-libhsail-rt: install-pdf-target-libhsail-rt
install-pdf-target-libhsail-rt: \
configure-target-libhsail-rt \
pdf-target-libhsail-rt
@: $(MAKE); $(unstage)
@[ -f $(TARGET_SUBDIR)/libhsail-rt/Makefile ] || exit 0; \
r=`${PWD_COMMAND}`; export r; \
s=`cd $(srcdir); ${PWD_COMMAND}`; export s; \
$(NORMAL_TARGET_EXPORTS) \
echo "Doing install-pdf in $(TARGET_SUBDIR)/libhsail-rt"; \
for flag in $(EXTRA_TARGET_FLAGS); do \
eval `echo "$$flag" | sed -e "s|^\([^=]*\)=\(.*\)|\1='\2'; export \1|"`; \
done; \
(cd $(TARGET_SUBDIR)/libhsail-rt && \
$(MAKE) $(BASE_FLAGS_TO_PASS) "AR=$${AR}" "AS=$${AS}" \
"CC=$${CC}" "CXX=$${CXX}" "LD=$${LD}" "NM=$${NM}" \
"RANLIB=$${RANLIB}" \
"DLLTOOL=$${DLLTOOL}" "WINDRES=$${WINDRES}" "WINDMC=$${WINDMC}" \
install-pdf) \
|| exit 1
@endif target-libhsail-rt
.PHONY: maybe-install-html-target-libhsail-rt install-html-target-libhsail-rt
maybe-install-html-target-libhsail-rt:
@if target-libhsail-rt
maybe-install-html-target-libhsail-rt: install-html-target-libhsail-rt
install-html-target-libhsail-rt: \
configure-target-libhsail-rt \
html-target-libhsail-rt
@: $(MAKE); $(unstage)
@[ -f $(TARGET_SUBDIR)/libhsail-rt/Makefile ] || exit 0; \
r=`${PWD_COMMAND}`; export r; \
s=`cd $(srcdir); ${PWD_COMMAND}`; export s; \
$(NORMAL_TARGET_EXPORTS) \
echo "Doing install-html in $(TARGET_SUBDIR)/libhsail-rt"; \
for flag in $(EXTRA_TARGET_FLAGS); do \
eval `echo "$$flag" | sed -e "s|^\([^=]*\)=\(.*\)|\1='\2'; export \1|"`; \
done; \
(cd $(TARGET_SUBDIR)/libhsail-rt && \
$(MAKE) $(BASE_FLAGS_TO_PASS) "AR=$${AR}" "AS=$${AS}" \
"CC=$${CC}" "CXX=$${CXX}" "LD=$${LD}" "NM=$${NM}" \
"RANLIB=$${RANLIB}" \
"DLLTOOL=$${DLLTOOL}" "WINDRES=$${WINDRES}" "WINDMC=$${WINDMC}" \
install-html) \
|| exit 1
@endif target-libhsail-rt
.PHONY: maybe-installcheck-target-libhsail-rt installcheck-target-libhsail-rt
maybe-installcheck-target-libhsail-rt:
@if target-libhsail-rt
maybe-installcheck-target-libhsail-rt: installcheck-target-libhsail-rt
installcheck-target-libhsail-rt: \
configure-target-libhsail-rt
@: $(MAKE); $(unstage)
@[ -f $(TARGET_SUBDIR)/libhsail-rt/Makefile ] || exit 0; \
r=`${PWD_COMMAND}`; export r; \
s=`cd $(srcdir); ${PWD_COMMAND}`; export s; \
$(NORMAL_TARGET_EXPORTS) \
echo "Doing installcheck in $(TARGET_SUBDIR)/libhsail-rt"; \
for flag in $(EXTRA_TARGET_FLAGS); do \
eval `echo "$$flag" | sed -e "s|^\([^=]*\)=\(.*\)|\1='\2'; export \1|"`; \
done; \
(cd $(TARGET_SUBDIR)/libhsail-rt && \
$(MAKE) $(BASE_FLAGS_TO_PASS) "AR=$${AR}" "AS=$${AS}" \
"CC=$${CC}" "CXX=$${CXX}" "LD=$${LD}" "NM=$${NM}" \
"RANLIB=$${RANLIB}" \
"DLLTOOL=$${DLLTOOL}" "WINDRES=$${WINDRES}" "WINDMC=$${WINDMC}" \
installcheck) \
|| exit 1
@endif target-libhsail-rt
.PHONY: maybe-mostlyclean-target-libhsail-rt mostlyclean-target-libhsail-rt
maybe-mostlyclean-target-libhsail-rt:
@if target-libhsail-rt
maybe-mostlyclean-target-libhsail-rt: mostlyclean-target-libhsail-rt
mostlyclean-target-libhsail-rt:
@: $(MAKE); $(unstage)
@[ -f $(TARGET_SUBDIR)/libhsail-rt/Makefile ] || exit 0; \
r=`${PWD_COMMAND}`; export r; \
s=`cd $(srcdir); ${PWD_COMMAND}`; export s; \
$(NORMAL_TARGET_EXPORTS) \
echo "Doing mostlyclean in $(TARGET_SUBDIR)/libhsail-rt"; \
for flag in $(EXTRA_TARGET_FLAGS); do \
eval `echo "$$flag" | sed -e "s|^\([^=]*\)=\(.*\)|\1='\2'; export \1|"`; \
done; \
(cd $(TARGET_SUBDIR)/libhsail-rt && \
$(MAKE) $(BASE_FLAGS_TO_PASS) "AR=$${AR}" "AS=$${AS}" \
"CC=$${CC}" "CXX=$${CXX}" "LD=$${LD}" "NM=$${NM}" \
"RANLIB=$${RANLIB}" \
"DLLTOOL=$${DLLTOOL}" "WINDRES=$${WINDRES}" "WINDMC=$${WINDMC}" \
mostlyclean) \
|| exit 1
@endif target-libhsail-rt
.PHONY: maybe-clean-target-libhsail-rt clean-target-libhsail-rt
maybe-clean-target-libhsail-rt:
@if target-libhsail-rt
maybe-clean-target-libhsail-rt: clean-target-libhsail-rt
clean-target-libhsail-rt:
@: $(MAKE); $(unstage)
@[ -f $(TARGET_SUBDIR)/libhsail-rt/Makefile ] || exit 0; \
r=`${PWD_COMMAND}`; export r; \
s=`cd $(srcdir); ${PWD_COMMAND}`; export s; \
$(NORMAL_TARGET_EXPORTS) \
echo "Doing clean in $(TARGET_SUBDIR)/libhsail-rt"; \
for flag in $(EXTRA_TARGET_FLAGS); do \
eval `echo "$$flag" | sed -e "s|^\([^=]*\)=\(.*\)|\1='\2'; export \1|"`; \
done; \
(cd $(TARGET_SUBDIR)/libhsail-rt && \
$(MAKE) $(BASE_FLAGS_TO_PASS) "AR=$${AR}" "AS=$${AS}" \
"CC=$${CC}" "CXX=$${CXX}" "LD=$${LD}" "NM=$${NM}" \
"RANLIB=$${RANLIB}" \
"DLLTOOL=$${DLLTOOL}" "WINDRES=$${WINDRES}" "WINDMC=$${WINDMC}" \
clean) \
|| exit 1
@endif target-libhsail-rt
.PHONY: maybe-distclean-target-libhsail-rt distclean-target-libhsail-rt
maybe-distclean-target-libhsail-rt:
@if target-libhsail-rt
maybe-distclean-target-libhsail-rt: distclean-target-libhsail-rt
distclean-target-libhsail-rt:
@: $(MAKE); $(unstage)
@[ -f $(TARGET_SUBDIR)/libhsail-rt/Makefile ] || exit 0; \
r=`${PWD_COMMAND}`; export r; \
s=`cd $(srcdir); ${PWD_COMMAND}`; export s; \
$(NORMAL_TARGET_EXPORTS) \
echo "Doing distclean in $(TARGET_SUBDIR)/libhsail-rt"; \
for flag in $(EXTRA_TARGET_FLAGS); do \
eval `echo "$$flag" | sed -e "s|^\([^=]*\)=\(.*\)|\1='\2'; export \1|"`; \
done; \
(cd $(TARGET_SUBDIR)/libhsail-rt && \
$(MAKE) $(BASE_FLAGS_TO_PASS) "AR=$${AR}" "AS=$${AS}" \
"CC=$${CC}" "CXX=$${CXX}" "LD=$${LD}" "NM=$${NM}" \
"RANLIB=$${RANLIB}" \
"DLLTOOL=$${DLLTOOL}" "WINDRES=$${WINDRES}" "WINDMC=$${WINDMC}" \
distclean) \
|| exit 1
@endif target-libhsail-rt
.PHONY: maybe-maintainer-clean-target-libhsail-rt maintainer-clean-target-libhsail-rt
maybe-maintainer-clean-target-libhsail-rt:
@if target-libhsail-rt
maybe-maintainer-clean-target-libhsail-rt: maintainer-clean-target-libhsail-rt
maintainer-clean-target-libhsail-rt:
@: $(MAKE); $(unstage)
@[ -f $(TARGET_SUBDIR)/libhsail-rt/Makefile ] || exit 0; \
r=`${PWD_COMMAND}`; export r; \
s=`cd $(srcdir); ${PWD_COMMAND}`; export s; \
$(NORMAL_TARGET_EXPORTS) \
echo "Doing maintainer-clean in $(TARGET_SUBDIR)/libhsail-rt"; \
for flag in $(EXTRA_TARGET_FLAGS); do \
eval `echo "$$flag" | sed -e "s|^\([^=]*\)=\(.*\)|\1='\2'; export \1|"`; \
done; \
(cd $(TARGET_SUBDIR)/libhsail-rt && \
$(MAKE) $(BASE_FLAGS_TO_PASS) "AR=$${AR}" "AS=$${AS}" \
"CC=$${CC}" "CXX=$${CXX}" "LD=$${LD}" "NM=$${NM}" \
"RANLIB=$${RANLIB}" \
"DLLTOOL=$${DLLTOOL}" "WINDRES=$${WINDRES}" "WINDMC=$${WINDMC}" \
maintainer-clean) \
|| exit 1
@endif target-libhsail-rt
.PHONY: configure-target-libphobos maybe-configure-target-libphobos
maybe-configure-target-libphobos:
@if gcc-bootstrap
@ -57045,14 +56569,6 @@ check-gcc-go:
(cd gcc && $(MAKE) $(GCC_FLAGS_TO_PASS) check-go);
check-go: check-gcc-go check-target-libgo check-gotools
.PHONY: check-gcc-brig check-brig
check-gcc-brig:
r=`${PWD_COMMAND}`; export r; \
s=`cd $(srcdir); ${PWD_COMMAND}`; export s; \
$(HOST_EXPORTS) \
(cd gcc && $(MAKE) $(GCC_FLAGS_TO_PASS) check-brig);
check-brig: check-gcc-brig check-target-libhsail-rt
.PHONY: check-gcc-d check-d
check-gcc-d:
r=`${PWD_COMMAND}`; export r; \
@ -60464,7 +59980,6 @@ configure-target-libquadmath: stage_last
configure-target-libgfortran: stage_last
configure-target-libobjc: stage_last
configure-target-libgo: stage_last
configure-target-libhsail-rt: stage_last
configure-target-libphobos: stage_last
configure-target-libtermcap: stage_last
configure-target-winsup: stage_last
@ -60499,7 +60014,6 @@ configure-target-libquadmath: maybe-all-gcc
configure-target-libgfortran: maybe-all-gcc
configure-target-libobjc: maybe-all-gcc
configure-target-libgo: maybe-all-gcc
configure-target-libhsail-rt: maybe-all-gcc
configure-target-libphobos: maybe-all-gcc
configure-target-libtermcap: maybe-all-gcc
configure-target-winsup: maybe-all-gcc
@ -61794,7 +61308,6 @@ configure-target-libquadmath: maybe-all-target-libgcc
configure-target-libgfortran: maybe-all-target-libgcc
configure-target-libobjc: maybe-all-target-libgcc
configure-target-libgo: maybe-all-target-libgcc
configure-target-libhsail-rt: maybe-all-target-libgcc
configure-target-libphobos: maybe-all-target-libgcc
configure-target-libtermcap: maybe-all-target-libgcc
configure-target-winsup: maybe-all-target-libgcc
@ -61832,8 +61345,6 @@ configure-target-libobjc: maybe-all-target-newlib maybe-all-target-libgloss
configure-target-libgo: maybe-all-target-newlib maybe-all-target-libgloss
configure-target-libhsail-rt: maybe-all-target-newlib maybe-all-target-libgloss
configure-target-libphobos: maybe-all-target-newlib maybe-all-target-libgloss
configure-target-libtermcap: maybe-all-target-newlib maybe-all-target-libgloss

24
configure vendored
View File

@ -2812,7 +2812,6 @@ target_libraries="target-libgcc \
target-newlib \
target-libgomp \
target-liboffloadmic \
target-libhsail-rt \
target-libatomic \
target-libitm \
target-libstdc++-v3 \
@ -3332,29 +3331,6 @@ $as_echo "yes" >&6; }
fi
fi
# Disable libhsail-rt on unsupported systems.
if test -d ${srcdir}/libhsail-rt; then
if test x$enable_libhsail_rt = x; then
{ $as_echo "$as_me:${as_lineno-$LINENO}: checking for libhsail-rt support" >&5
$as_echo_n "checking for libhsail-rt support... " >&6; }
if (srcdir=${srcdir}/libhsail-rt; \
. ${srcdir}/configure.tgt; \
test -n "$UNSUPPORTED")
then
{ $as_echo "$as_me:${as_lineno-$LINENO}: result: no" >&5
$as_echo "no" >&6; }
unsupported_languages="$unsupported_languages brig"
# This implicitly disables also target-libhsail-rt as it won't
# get added to the build without BRIG FE.
else
{ $as_echo "$as_me:${as_lineno-$LINENO}: result: yes" >&5
$as_echo "yes" >&6; }
fi
fi
fi
# Disable libquadmath for some systems.
case "${target}" in
avr-*-*)

21
configure.ac
View File

@ -151,7 +151,6 @@ target_libraries="target-libgcc \
target-newlib \
target-libgomp \
target-liboffloadmic \
target-libhsail-rt \
target-libatomic \
target-libitm \
target-libstdc++-v3 \
@ -609,26 +608,6 @@ if test -d ${srcdir}/libvtv; then
fi
fi
# Disable libhsail-rt on unsupported systems.
if test -d ${srcdir}/libhsail-rt; then
if test x$enable_libhsail_rt = x; then
AC_MSG_CHECKING([for libhsail-rt support])
if (srcdir=${srcdir}/libhsail-rt; \
. ${srcdir}/configure.tgt; \
test -n "$UNSUPPORTED")
then
AC_MSG_RESULT([no])
unsupported_languages="$unsupported_languages brig"
# This implicitly disables also target-libhsail-rt as it won't
# get added to the build without BRIG FE.
else
AC_MSG_RESULT([yes])
fi
fi
fi
# Disable libquadmath for some systems.
case "${target}" in
avr-*-*)

4
contrib/gcc_update
View File

@ -141,10 +141,6 @@ libgomp/testsuite/Makefile.in: libgomp/testsuite/Makefile.am libgomp/aclocal.m4
libgomp/configure.ac: libgomp/plugin/configfrag.ac
libgomp/configure: libgomp/configure.ac libgomp/aclocal.m4
libgomp/config.h.in: libgomp/configure.ac libgomp/aclocal.m4
libhsail-rt/aclocal.m4: libhsail-rt/configure.ac
libhsail-rt/Makefile.in: libhsail-rt/Makefile.am libhsail-rt/aclocal.m4
libhsail-rt/configure: libhsail-rt/configure.ac libhsail-rt/aclocal.m4
libhsail-rt/target-config.h.in: libhsail-rt/configure.ac libhsail-rt/aclocal.m4
libitm/aclocal.m4: libitm/configure.ac libitm/acinclude.m4
libitm/Makefile.in: libitm/Makefile.am libitm/aclocal.m4
libitm/testsuite/Makefile.in: libitm/testsuite/Makefile.am libitm/aclocal.m4

2
contrib/update-copyright.py
View File

@ -755,7 +755,6 @@ class GCCCmdLine (CmdLine):
self.add_dir ('libgfortran')
# libgo is imported from upstream.
self.add_dir ('libgomp')
self.add_dir ('libhsail-rt')
self.add_dir ('libiberty')
self.add_dir ('libitm')
self.add_dir ('libobjc')
@ -782,7 +781,6 @@ class GCCCmdLine (CmdLine):
'libgcc',
'libgfortran',
'libgomp',
'libhsail-rt',
'libiberty',
'libitm',
'libobjc',

164
libhsail-rt/ChangeLog
View File

@ -1,164 +0,0 @@
2021-01-05 Samuel Thibault <samuel.thibault@ens-lyon.org>
* configure: Re-generate.
2020-12-05 Iain Sandoe <iain@sandoe.co.uk>
PR target/97865
* configure: Regenerate.
2020-11-29 John David Anglin <danglin@gcc.gnu.org>
* configure: Regenerate.
2020-01-24 Maciej W. Rozycki <macro@wdc.com>
* configure.ac: Handle `--with-toolexeclibdir='.
* Makefile.in: Regenerate.
* aclocal.m4: Regenerate.
* configure: Regenerate.
2020-01-01 Jakub Jelinek <jakub@redhat.com>
Update copyright years.
2019-09-27 Maciej W. Rozycki <macro@wdc.com>
* configure: Regenerate.
2019-01-01 Jakub Jelinek <jakub@redhat.com>
Update copyright years.
2018-10-31 Joseph Myers <joseph@codesourcery.com>
PR bootstrap/82856
* configure.ac: Remove AC_PREREQ.
* Makefile.in, aclocal.m4, configure: Regenerate.
2018-05-04 Pekka Jääskeläinen <pekka.jaaskelainen@parmance.com>
* include/internal/phsa-rt.h: Whitespace cleanup.
* include/internal/workitems.h: Store work item ID data to easily
accessible locations.
* rt/workitems.c: Same.
2018-05-04 Pekka Jääskeläinen <pekka.jaaskelainen@parmance.com>
* rt/workitems.c: Fix an alloca stack underflow.
2018-04-18 David Malcolm <dmalcolm@redhat.com>
PR jit/85384
* configure: Regenerate.
2018-01-03 Jakub Jelinek <jakub@redhat.com>
Update copyright years.
2017-09-27 Pekka Jääskeläinen <pekka.jaaskelainen@parmance.com>
* include/internal/phsa-rt.h: Support for improved group segment
handling with a stack-like allocation scheme.
* include/internal/workitems.h: Likewise.
* rt/workitems.c: Likewise.
2017-09-25 Pekka Jääskeläinen <pekka.jaaskelainen@parmance.com>
* rt/workitems.c: Assume the host runtime allocates the work group
memory.
2017-05-03 Pekka Jääskeläinen <pekka.jaaskelainen@parmance.com>
* rt/workitems.c: Removed a leftover comment.
* rt/arithmetic.c (__hsail_class_f32, __hsail_class_f64): Fix the
check for signaling/non-signalling NaN. Add class_f64 default
implementation.
2017-02-01 Jakub Jelinek <jakub@redhat.com>
* configure.tgt: Fix i?86-*-linux* entry.
* rt/sat_arithmetic.c (__hsail_sat_add_u32, __hsail_sat_add_u64,
__hsail_sat_add_s32, __hsail_sat_add_s64): Use __builtin_add_overflow.
(__hsail_sat_sub_u8, __hsail_sat_sub_u16): Remove pointless for overflow
over maximum.
(__hsail_sat_sub_u32, __hsail_sat_sub_u64, __hsail_sat_sub_s32,
__hsail_sat_sub_s64): Use __builtin_sub_overflow.
(__hsail_sat_mul_u32, __hsail_sat_mul_u64, __hsail_sat_mul_s32,
__hsail_sat_mul_s64): Use __builtin_mul_overflow.
* rt/arithmetic.c (__hsail_borrow_u32, __hsail_borrow_u64): Use
__builtin_sub_overflow_p.
(__hsail_carry_u32, __hsail_carry_u64): Use __builtin_add_overflow_p.
* rt/misc.c (__hsail_groupbaseptr, __hsail_kernargbaseptr_u64):
Cast pointers to uintptr_t first before casting to some other integral
type.
* rt/segment.c (__hsail_segmentp_private, __hsail_segmentp_group): Likewise.
* rt/queue.c (__hsail_ldqueuereadindex, __hsail_ldqueuewriteindex,
__hsail_addqueuewriteindex, __hsail_casqueuewriteindex,
__hsail_stqueuereadindex, __hsail_stqueuewriteindex): Cast integral value
to uintptr_t first before casting to pointer.
* rt/workitems.c (__hsail_alloca_pop_frame): Cast memcpy first argument to
void * to avoid warning.
2017-01-27 Pekka Jääskeläinen <pekka.jaaskelainen@parmance.com>
* configure.tgt: Moved the white list of supported targets here
from configure.ac. Added i[3456789]86-*-linux* as a supported env
for the BRIG FE.
* README: Added a proper description of what libhsail-rt is.
2017-01-26 Jakub Jelinek <jakub@redhat.com>
Update copyright years.
2017-01-25 Thomas Schwinge <thomas@codesourcery.com>
* config.h.in: Remove stale file.
* configure.ac: Don't instantiate AC_CONFIG_MACRO_DIR.
* configure: Regenerate.
2017-01-25 Jakub Jelinek <jakub@redhat.com>
PR other/79046
* configure.ac: Add GCC_BASE_VER.
* Makefile.am (gcc_version): Use @get_gcc_base_ver@ instead of cat to
get version from BASE-VER file.
(ACLOCAL_AMFLAGS): Set to -I .. -I ../config .
* aclocal.m4: Regenerated.
* configure: Regenerated.
* Makefile.in: Regenerated.
2017-01-24 Pekka Jääskeläinen <pekka@parmance.com>
Martin Jambor <mjambor@suse.cz>
* Makefile.am: New file.
* target-config.h.in: Likewise.
* configure.ac: Likewise.
* configure: Likewise.
* config.h.in: Likewise.
* aclocal.m4: Likewise.
* README: Likewise.
* Makefile.in: Likewise.
* include/internal/fibers.h: Likewise.
* include/internal/phsa-queue-interface.h: Likewise.
* include/internal/phsa-rt.h: Likewise.
* include/internal/workitems.h: Likewise.
* rt/arithmetic.c: Likewise.
* rt/atomics.c: Likewise.
* rt/bitstring.c: Likewise.
* rt/fbarrier.c: Likewise.
* rt/fibers.c: Likewise.
* rt/fp16.c: Likewise.
* rt/misc.c: Likewise.
* rt/multimedia.c: Likewise.
* rt/queue.c: Likewise.
* rt/sat_arithmetic.c: Likewise.
* rt/segment.c: Likewise.
* rt/workitems.c: Likewise.
Copyright (C) 2017-2021 Free Software Foundation, Inc.
Copying and distribution of this file, with or without modification,
are permitted in any medium without royalty provided the copyright
notice and this notice are preserved.

122
libhsail-rt/Makefile.am
View File

@ -1,122 +0,0 @@
# Makefile.am -- libhsail-rt library Makefile.
# Starting point copied from libcilkrts:
# @copyright
# Copyright (C) 2011, 2013, Intel Corporation
# All rights reserved.
#
# @copyright
# Redistribution and use in source and binary forms, with or without
# modification, are permitted provided that the following conditions
# are met:
#
# * Redistributions of source code must retain the above copyright
# notice, this list of conditions and the following disclaimer.
# * Redistributions in binary form must reproduce the above copyright
# notice, this list of conditions and the following disclaimer in
# the documentation and/or other materials provided with the
# distribution.
# * Neither the name of Intel Corporation nor the names of its
# contributors may be used to endorse or promote products derived
# from this software without specific prior written permission.
#
# @copyright
# THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS
# "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT
# LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR
# A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT
# HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT,
# INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
# BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS
# OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED
# AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT
# LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY
# WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE
# POSSIBILITY OF SUCH DAMAGE.
# libhsail-rt modifications:
# Contributed by Pekka Jaaskelainen <pekka.jaaskelainen@parmance.com>
# for General Processor Tech.
# Use of this source code is governed by a BSD-style
# license that can be found in the LICENSE file.
# Process this file with autoreconf to produce Makefile.in.
AUTOMAKE_OPTIONS = foreign subdir-objects
gcc_version := $(shell @get_gcc_base_ver@ $(top_srcdir)/../gcc/BASE-VER)
MAINT_CHARSET = latin1
mkinstalldirs = $(SHELL) $(toplevel_srcdir)/mkinstalldirs
ACLOCAL_AMFLAGS = -I .. -I ../config
WARN_CFLAGS = $(WARN_FLAGS) $(WERROR)
# -I/-D flags to pass when compiling.
AM_CPPFLAGS = -I$(srcdir)/rt -I$(srcdir)/include/internal
AM_CFLAGS = \
-I $(srcdir)/../include \
-I $(srcdir)/../libgcc \
-I $(MULTIBUILDTOP)../../gcc/include $(PTH_CFLAGS)
toolexeclib_LTLIBRARIES = libhsail-rt.la
runtime_files = \
rt/arithmetic.c \
rt/atomics.c \
rt/bitstring.c \
rt/fbarrier.c \
rt/fp16.c \
rt/misc.c \
rt/multimedia.c \
rt/queue.c \
rt/sat_arithmetic.c \
rt/segment.c \
rt/workitems.c \
rt/fibers.c
libhsail_rt_la_SOURCES = $(runtime_files)
libhsail_rt_la_LDFLAGS = -rpath '$(libdir)'
# Work around what appears to be a GNU make bug handling MAKEFLAGS
# values defined in terms of make variables, as is the case for CC and
# friends when we are called from the top level Makefile.
AM_MAKEFLAGS = \
"AR_FLAGS=$(AR_FLAGS)" \
"CC_FOR_BUILD=$(CC_FOR_BUILD)" \
"CFLAGS=$(CFLAGS)" \
"CXXFLAGS=$(CXXFLAGS)" \
"CFLAGS_FOR_BUILD=$(CFLAGS_FOR_BUILD)" \
"CFLAGS_FOR_TARGET=$(CFLAGS_FOR_TARGET)" \
"INSTALL=$(INSTALL)" \
"INSTALL_DATA=$(INSTALL_DATA)" \
"INSTALL_PROGRAM=$(INSTALL_PROGRAM)" \
"INSTALL_SCRIPT=$(INSTALL_SCRIPT)" \
"JC1FLAGS=$(JC1FLAGS)" \
"LDFLAGS=$(LDFLAGS)" \
"LIBCFLAGS=$(LIBCFLAGS)" \
"LIBCFLAGS_FOR_TARGET=$(LIBCFLAGS_FOR_TARGET)" \
"MAKE=$(MAKE)" \
"MAKEINFO=$(MAKEINFO) $(MAKEINFOFLAGS)" \
"PICFLAG=$(PICFLAG)" \
"PICFLAG_FOR_TARGET=$(PICFLAG_FOR_TARGET)" \
"SHELL=$(SHELL)" \
"RUNTESTFLAGS=$(RUNTESTFLAGS)" \
"exec_prefix=$(exec_prefix)" \
"infodir=$(infodir)" \
"libdir=$(libdir)" \
"prefix=$(prefix)" \
"includedir=$(includedir)" \
"AR=$(AR)" \
"AS=$(AS)" \
"LD=$(LD)" \
"LIBCFLAGS=$(LIBCFLAGS)" \
"NM=$(NM)" \
"PICFLAG=$(PICFLAG)" \
"RANLIB=$(RANLIB)" \
"DESTDIR=$(DESTDIR)"
MAKEOVERRIDES=

817
libhsail-rt/Makefile.in
View File

@ -1,817 +0,0 @@
# Makefile.in generated by automake 1.15.1 from Makefile.am.
# @configure_input@
# Copyright (C) 1994-2017 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.
# This program is distributed in the hope that it will be useful,
# but WITHOUT ANY WARRANTY, to the extent permitted by law; without
# even the implied warranty of MERCHANTABILITY or FITNESS FOR A
# PARTICULAR PURPOSE.
@SET_MAKE@
# Makefile.am -- libhsail-rt library Makefile.
# Starting point copied from libcilkrts:
# @copyright
# Copyright (C) 2011, 2013, Intel Corporation
# All rights reserved.
#
# @copyright
# Redistribution and use in source and binary forms, with or without
# modification, are permitted provided that the following conditions
# are met:
#
# * Redistributions of source code must retain the above copyright
# notice, this list of conditions and the following disclaimer.
# * Redistributions in binary form must reproduce the above copyright
# notice, this list of conditions and the following disclaimer in
# the documentation and/or other materials provided with the
# distribution.
# * Neither the name of Intel Corporation nor the names of its
# contributors may be used to endorse or promote products derived
# from this software without specific prior written permission.
#
# @copyright
# THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS
# "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT
# LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR
# A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT
# HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT,
# INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
# BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS
# OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED
# AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT
# LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY
# WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE
# POSSIBILITY OF SUCH DAMAGE.
# libhsail-rt modifications:
# Contributed by Pekka Jaaskelainen <pekka.jaaskelainen@parmance.com>
# for General Processor Tech.
# Use of this source code is governed by a BSD-style
# license that can be found in the LICENSE file.
# Process this file with autoreconf to produce Makefile.in.
VPATH = @srcdir@
am__is_gnu_make = { \
if test -z '$(MAKELEVEL)'; then \
false; \
elif test -n '$(MAKE_HOST)'; then \
true; \
elif test -n '$(MAKE_VERSION)' && test -n '$(CURDIR)'; then \
true; \
else \
false; \
fi; \
}
am__make_running_with_option = \
case $${target_option-} in \
?) ;; \
*) echo "am__make_running_with_option: internal error: invalid" \
"target option '$${target_option-}' specified" >&2; \
exit 1;; \
esac; \
has_opt=no; \
sane_makeflags=$$MAKEFLAGS; \
if $(am__is_gnu_make); then \
sane_makeflags=$$MFLAGS; \
else \
case $$MAKEFLAGS in \
*\\[\ \ ]*) \
bs=\\; \
sane_makeflags=`printf '%s\n' "$$MAKEFLAGS" \
| sed "s/$$bs$$bs[$$bs $$bs ]*//g"`;; \
esac; \
fi; \
skip_next=no; \
strip_trailopt () \
{ \
flg=`printf '%s\n' "$$flg" | sed "s/$$1.*$$//"`; \
}; \
for flg in $$sane_makeflags; do \
test $$skip_next = yes && { skip_next=no; continue; }; \
case $$flg in \
*=*|--*) continue;; \
-*I) strip_trailopt 'I'; skip_next=yes;; \
-*I?*) strip_trailopt 'I';; \
-*O) strip_trailopt 'O'; skip_next=yes;; \
-*O?*) strip_trailopt 'O';; \
-*l) strip_trailopt 'l'; skip_next=yes;; \
-*l?*) strip_trailopt 'l';; \
-[dEDm]) skip_next=yes;; \
-[JT]) skip_next=yes;; \
esac; \
case $$flg in \
*$$target_option*) has_opt=yes; break;; \
esac; \
done; \
test $$has_opt = yes
am__make_dryrun = (target_option=n; $(am__make_running_with_option))
am__make_keepgoing = (target_option=k; $(am__make_running_with_option))
pkgdatadir = $(datadir)/@PACKAGE@
pkgincludedir = $(includedir)/@PACKAGE@
pkglibdir = $(libdir)/@PACKAGE@
pkglibexecdir = $(libexecdir)/@PACKAGE@
am__cd = CDPATH="$${ZSH_VERSION+.}$(PATH_SEPARATOR)" && cd
install_sh_DATA = $(install_sh) -c -m 644
install_sh_PROGRAM = $(install_sh) -c
install_sh_SCRIPT = $(install_sh) -c
INSTALL_HEADER = $(INSTALL_DATA)
transform = $(program_transform_name)
NORMAL_INSTALL = :
PRE_INSTALL = :
POST_INSTALL = :
NORMAL_UNINSTALL = :
PRE_UNINSTALL = :
POST_UNINSTALL = :
build_triplet = @build@
host_triplet = @host@
target_triplet = @target@
subdir = .
ACLOCAL_M4 = $(top_srcdir)/aclocal.m4
am__aclocal_m4_deps = $(top_srcdir)/../config/acx.m4 \
$(top_srcdir)/../config/depstand.m4 \
$(top_srcdir)/../config/lead-dot.m4 \
$(top_srcdir)/../config/override.m4 \
$(top_srcdir)/../config/toolexeclibdir.m4 \
$(top_srcdir)/../libtool.m4 $(top_srcdir)/../ltoptions.m4 \
$(top_srcdir)/../ltsugar.m4 $(top_srcdir)/../ltversion.m4 \
$(top_srcdir)/../lt~obsolete.m4 $(top_srcdir)/configure.ac
am__configure_deps = $(am__aclocal_m4_deps) $(CONFIGURE_DEPENDENCIES) \
$(ACLOCAL_M4)
DIST_COMMON = $(srcdir)/Makefile.am $(top_srcdir)/configure \
$(am__configure_deps)
am__CONFIG_DISTCLEAN_FILES = config.status config.cache config.log \
configure.lineno config.status.lineno
CONFIG_HEADER = target-config.h
CONFIG_CLEAN_FILES =
CONFIG_CLEAN_VPATH_FILES =
am__vpath_adj_setup = srcdirstrip=`echo "$(srcdir)" | sed 's|.|.|g'`;
am__vpath_adj = case $$p in \
$(srcdir)/*) f=`echo "$$p" | sed "s|^$$srcdirstrip/||"`;; \
*) f=$$p;; \
esac;
am__strip_dir = f=`echo $$p | sed -e 's|^.*/||'`;
am__install_max = 40
am__nobase_strip_setup = \
srcdirstrip=`echo "$(srcdir)" | sed 's/[].[^$$\\*|]/\\\\&/g'`
am__nobase_strip = \
for p in $$list; do echo "$$p"; done | sed -e "s|$$srcdirstrip/||"
am__nobase_list = $(am__nobase_strip_setup); \
for p in $$list; do echo "$$p $$p"; done | \
sed "s| $$srcdirstrip/| |;"' / .*\//!s/ .*/ ./; s,\( .*\)/[^/]*$$,\1,' | \
$(AWK) 'BEGIN { files["."] = "" } { files[$$2] = files[$$2] " " $$1; \
if (++n[$$2] == $(am__install_max)) \
{ print $$2, files[$$2]; n[$$2] = 0; files[$$2] = "" } } \
END { for (dir in files) print dir, files[dir] }'
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)"
LTLIBRARIES = $(toolexeclib_LTLIBRARIES)
libhsail_rt_la_LIBADD =
am__dirstamp = $(am__leading_dot)dirstamp
am__objects_1 = rt/arithmetic.lo rt/atomics.lo rt/bitstring.lo \
rt/fbarrier.lo rt/fp16.lo rt/misc.lo rt/multimedia.lo \
rt/queue.lo rt/sat_arithmetic.lo rt/segment.lo rt/workitems.lo \
rt/fibers.lo
am_libhsail_rt_la_OBJECTS = $(am__objects_1)
libhsail_rt_la_OBJECTS = $(am_libhsail_rt_la_OBJECTS)
AM_V_lt = $(am__v_lt_@AM_V@)
am__v_lt_ = $(am__v_lt_@AM_DEFAULT_V@)
am__v_lt_0 = --silent
am__v_lt_1 =
libhsail_rt_la_LINK = $(LIBTOOL) $(AM_V_lt) --tag=CC \
$(AM_LIBTOOLFLAGS) $(LIBTOOLFLAGS) --mode=link $(CCLD) \
$(AM_CFLAGS) $(CFLAGS) $(libhsail_rt_la_LDFLAGS) $(LDFLAGS) -o \
$@
AM_V_P = $(am__v_P_@AM_V@)
am__v_P_ = $(am__v_P_@AM_DEFAULT_V@)
am__v_P_0 = false
am__v_P_1 = :
AM_V_GEN = $(am__v_GEN_@AM_V@)
am__v_GEN_ = $(am__v_GEN_@AM_DEFAULT_V@)
am__v_GEN_0 = @echo " GEN " $@;
am__v_GEN_1 =
AM_V_at = $(am__v_at_@AM_V@)
am__v_at_ = $(am__v_at_@AM_DEFAULT_V@)
am__v_at_0 = @
am__v_at_1 =
DEFAULT_INCLUDES = -I.@am__isrc@
depcomp = $(SHELL) $(top_srcdir)/../depcomp
am__depfiles_maybe = depfiles
am__mv = mv -f
COMPILE = $(CC) $(DEFS) $(DEFAULT_INCLUDES) $(INCLUDES) $(AM_CPPFLAGS) \
$(CPPFLAGS) $(AM_CFLAGS) $(CFLAGS)
LTCOMPILE = $(LIBTOOL) $(AM_V_lt) --tag=CC $(AM_LIBTOOLFLAGS) \
$(LIBTOOLFLAGS) --mode=compile $(CC) $(DEFS) \
$(DEFAULT_INCLUDES) $(INCLUDES) $(AM_CPPFLAGS) $(CPPFLAGS) \
$(AM_CFLAGS) $(CFLAGS)
AM_V_CC = $(am__v_CC_@AM_V@)
am__v_CC_ = $(am__v_CC_@AM_DEFAULT_V@)
am__v_CC_0 = @echo " CC " $@;
am__v_CC_1 =
CCLD = $(CC)
LINK = $(LIBTOOL) $(AM_V_lt) --tag=CC $(AM_LIBTOOLFLAGS) \
$(LIBTOOLFLAGS) --mode=link $(CCLD) $(AM_CFLAGS) $(CFLAGS) \
$(AM_LDFLAGS) $(LDFLAGS) -o $@
AM_V_CCLD = $(am__v_CCLD_@AM_V@)
am__v_CCLD_ = $(am__v_CCLD_@AM_DEFAULT_V@)
am__v_CCLD_0 = @echo " CCLD " $@;
am__v_CCLD_1 =
SOURCES = $(libhsail_rt_la_SOURCES)
am__can_run_installinfo = \
case $$AM_UPDATE_INFO_DIR in \
n|no|NO) false;; \
*) (install-info --version) >/dev/null 2>&1;; \
esac
am__tagged_files = $(HEADERS) $(SOURCES) $(TAGS_FILES) \
$(LISP)target-config.h.in
# Read a list of newline-separated strings from the standard input,
# and print each of them once, without duplicates. Input order is
# *not* preserved.
am__uniquify_input = $(AWK) '\
BEGIN { nonempty = 0; } \
{ items[$$0] = 1; nonempty = 1; } \
END { if (nonempty) { for (i in items) print i; }; } \
'
# Make sure the list of sources is unique. This is necessary because,
# e.g., the same source file might be shared among _SOURCES variables
# for different programs/libraries.
am__define_uniq_tagged_files = \
list='$(am__tagged_files)'; \
unique=`for i in $$list; do \
if test -f "$$i"; then echo $$i; else echo $(srcdir)/$$i; fi; \
done | $(am__uniquify_input)`
ETAGS = etags
CTAGS = ctags
CSCOPE = cscope
AM_RECURSIVE_TARGETS = cscope
ACLOCAL = @ACLOCAL@
AMTAR = @AMTAR@
AM_DEFAULT_VERBOSITY = @AM_DEFAULT_VERBOSITY@
AR = @AR@
AUTOCONF = @AUTOCONF@
AUTOHEADER = @AUTOHEADER@
AUTOMAKE = @AUTOMAKE@
AWK = @AWK@
CC = @CC@
CCDEPMODE = @CCDEPMODE@
CFLAGS = @CFLAGS@
CPP = @CPP@
CPPFLAGS = @CPPFLAGS@
CXX = @CXX@
CXXCPP = @CXXCPP@
CXXDEPMODE = @CXXDEPMODE@
CXXFLAGS = @CXXFLAGS@
CYGPATH_W = @CYGPATH_W@
DEFS = @DEFS@
DEPDIR = @DEPDIR@
DSYMUTIL = @DSYMUTIL@
DUMPBIN = @DUMPBIN@
ECHO_C = @ECHO_C@
ECHO_N = @ECHO_N@
ECHO_T = @ECHO_T@
EGREP = @EGREP@
EXEEXT = @EXEEXT@
FGREP = @FGREP@
GREP = @GREP@
INSTALL = @INSTALL@
INSTALL_DATA = @INSTALL_DATA@
INSTALL_PROGRAM = @INSTALL_PROGRAM@
INSTALL_SCRIPT = @INSTALL_SCRIPT@
INSTALL_STRIP_PROGRAM = @INSTALL_STRIP_PROGRAM@
LD = @LD@
LDFLAGS = @LDFLAGS@
LIBOBJS = @LIBOBJS@
LIBS = @LIBS@
LIBTOOL = @LIBTOOL@
LIPO = @LIPO@
LN_S = @LN_S@
LTLIBOBJS = @LTLIBOBJS@
MAINT = @MAINT@
MAKEINFO = @MAKEINFO@
MKDIR_P = @MKDIR_P@
NM = @NM@
NMEDIT = @NMEDIT@
OBJDUMP = @OBJDUMP@
OBJEXT = @OBJEXT@
OTOOL = @OTOOL@
OTOOL64 = @OTOOL64@
PACKAGE = @PACKAGE@
PACKAGE_BUGREPORT = @PACKAGE_BUGREPORT@
PACKAGE_NAME = @PACKAGE_NAME@
PACKAGE_STRING = @PACKAGE_STRING@
PACKAGE_TARNAME = @PACKAGE_TARNAME@
PACKAGE_URL = @PACKAGE_URL@
PACKAGE_VERSION = @PACKAGE_VERSION@
PATH_SEPARATOR = @PATH_SEPARATOR@
RANLIB = @RANLIB@
SED = @SED@
SET_MAKE = @SET_MAKE@
SHELL = @SHELL@
STRIP = @STRIP@
VERSION = @VERSION@
abs_builddir = @abs_builddir@
abs_srcdir = @abs_srcdir@
abs_top_builddir = @abs_top_builddir@
abs_top_srcdir = @abs_top_srcdir@
ac_ct_CC = @ac_ct_CC@
ac_ct_CXX = @ac_ct_CXX@
ac_ct_DUMPBIN = @ac_ct_DUMPBIN@
am__include = @am__include@
am__leading_dot = @am__leading_dot@
am__quote = @am__quote@
am__tar = @am__tar@
am__untar = @am__untar@
bindir = @bindir@
build = @build@
build_alias = @build_alias@
build_cpu = @build_cpu@
build_os = @build_os@
build_vendor = @build_vendor@
builddir = @builddir@
config_dir = @config_dir@
datadir = @datadir@
datarootdir = @datarootdir@
docdir = @docdir@
dvidir = @dvidir@
exec_prefix = @exec_prefix@
get_gcc_base_ver = @get_gcc_base_ver@
host = @host@
host_alias = @host_alias@
host_cpu = @host_cpu@
host_os = @host_os@
host_vendor = @host_vendor@
htmldir = @htmldir@
includedir = @includedir@
infodir = @infodir@
install_sh = @install_sh@
libdir = @libdir@
libexecdir = @libexecdir@
localedir = @localedir@
localstatedir = @localstatedir@
mandir = @mandir@
mkdir_p = @mkdir_p@
oldincludedir = @oldincludedir@
pdfdir = @pdfdir@
prefix = @prefix@
program_transform_name = @program_transform_name@
psdir = @psdir@
sbindir = @sbindir@
sharedstatedir = @sharedstatedir@
srcdir = @srcdir@
sysconfdir = @sysconfdir@
target = @target@
target_alias = @target_alias@
target_cpu = @target_cpu@
target_os = @target_os@
target_vendor = @target_vendor@
toolexecdir = @toolexecdir@
toolexeclibdir = @toolexeclibdir@
top_build_prefix = @top_build_prefix@
top_builddir = @top_builddir@
top_srcdir = @top_srcdir@
AUTOMAKE_OPTIONS = foreign subdir-objects
gcc_version := $(shell @get_gcc_base_ver@ $(top_srcdir)/../gcc/BASE-VER)
MAINT_CHARSET = latin1
mkinstalldirs = $(SHELL) $(toplevel_srcdir)/mkinstalldirs
ACLOCAL_AMFLAGS = -I .. -I ../config
WARN_CFLAGS = $(WARN_FLAGS) $(WERROR)
# -I/-D flags to pass when compiling.
AM_CPPFLAGS = -I$(srcdir)/rt -I$(srcdir)/include/internal
AM_CFLAGS = \
-I $(srcdir)/../include \
-I $(srcdir)/../libgcc \
-I $(MULTIBUILDTOP)../../gcc/include $(PTH_CFLAGS)
toolexeclib_LTLIBRARIES = libhsail-rt.la
runtime_files = \
rt/arithmetic.c \
rt/atomics.c \
rt/bitstring.c \
rt/fbarrier.c \
rt/fp16.c \
rt/misc.c \
rt/multimedia.c \
rt/queue.c \
rt/sat_arithmetic.c \
rt/segment.c \
rt/workitems.c \
rt/fibers.c
libhsail_rt_la_SOURCES = $(runtime_files)
libhsail_rt_la_LDFLAGS = -rpath '$(libdir)'
# Work around what appears to be a GNU make bug handling MAKEFLAGS
# values defined in terms of make variables, as is the case for CC and
# friends when we are called from the top level Makefile.
AM_MAKEFLAGS = \
"AR_FLAGS=$(AR_FLAGS)" \
"CC_FOR_BUILD=$(CC_FOR_BUILD)" \
"CFLAGS=$(CFLAGS)" \
"CXXFLAGS=$(CXXFLAGS)" \
"CFLAGS_FOR_BUILD=$(CFLAGS_FOR_BUILD)" \
"CFLAGS_FOR_TARGET=$(CFLAGS_FOR_TARGET)" \
"INSTALL=$(INSTALL)" \
"INSTALL_DATA=$(INSTALL_DATA)" \
"INSTALL_PROGRAM=$(INSTALL_PROGRAM)" \
"INSTALL_SCRIPT=$(INSTALL_SCRIPT)" \
"JC1FLAGS=$(JC1FLAGS)" \
"LDFLAGS=$(LDFLAGS)" \
"LIBCFLAGS=$(LIBCFLAGS)" \
"LIBCFLAGS_FOR_TARGET=$(LIBCFLAGS_FOR_TARGET)" \
"MAKE=$(MAKE)" \
"MAKEINFO=$(MAKEINFO) $(MAKEINFOFLAGS)" \
"PICFLAG=$(PICFLAG)" \
"PICFLAG_FOR_TARGET=$(PICFLAG_FOR_TARGET)" \
"SHELL=$(SHELL)" \
"RUNTESTFLAGS=$(RUNTESTFLAGS)" \
"exec_prefix=$(exec_prefix)" \
"infodir=$(infodir)" \
"libdir=$(libdir)" \
"prefix=$(prefix)" \
"includedir=$(includedir)" \
"AR=$(AR)" \
"AS=$(AS)" \
"LD=$(LD)" \
"LIBCFLAGS=$(LIBCFLAGS)" \
"NM=$(NM)" \
"PICFLAG=$(PICFLAG)" \
"RANLIB=$(RANLIB)" \
"DESTDIR=$(DESTDIR)"
MAKEOVERRIDES =
all: target-config.h
$(MAKE) $(AM_MAKEFLAGS) all-am
.SUFFIXES:
.SUFFIXES: .c .lo .o .obj
am--refresh: Makefile
@:
$(srcdir)/Makefile.in: @MAINTAINER_MODE_TRUE@ $(srcdir)/Makefile.am $(am__configure_deps)
@for dep in $?; do \
case '$(am__configure_deps)' in \
*$$dep*) \
echo ' cd $(srcdir) && $(AUTOMAKE) --foreign'; \
$(am__cd) $(srcdir) && $(AUTOMAKE) --foreign \
&& exit 0; \
exit 1;; \
esac; \
done; \
echo ' cd $(top_srcdir) && $(AUTOMAKE) --foreign Makefile'; \
$(am__cd) $(top_srcdir) && \
$(AUTOMAKE) --foreign Makefile
Makefile: $(srcdir)/Makefile.in $(top_builddir)/config.status
@case '$?' in \
*config.status*) \
echo ' $(SHELL) ./config.status'; \
$(SHELL) ./config.status;; \
*) \
echo ' cd $(top_builddir) && $(SHELL) ./config.status $@ $(am__depfiles_maybe)'; \
cd $(top_builddir) && $(SHELL) ./config.status $@ $(am__depfiles_maybe);; \
esac;
$(top_builddir)/config.status: $(top_srcdir)/configure $(CONFIG_STATUS_DEPENDENCIES)
$(SHELL) ./config.status --recheck
$(top_srcdir)/configure: @MAINTAINER_MODE_TRUE@ $(am__configure_deps)
$(am__cd) $(srcdir) && $(AUTOCONF)
$(ACLOCAL_M4): @MAINTAINER_MODE_TRUE@ $(am__aclocal_m4_deps)
$(am__cd) $(srcdir) && $(ACLOCAL) $(ACLOCAL_AMFLAGS)
$(am__aclocal_m4_deps):
target-config.h: stamp-h1
@test -f $@ || rm -f stamp-h1
@test -f $@ || $(MAKE) $(AM_MAKEFLAGS) stamp-h1
stamp-h1: $(srcdir)/target-config.h.in $(top_builddir)/config.status
@rm -f stamp-h1
cd $(top_builddir) && $(SHELL) ./config.status target-config.h
$(srcdir)/target-config.h.in: @MAINTAINER_MODE_TRUE@ $(am__configure_deps)
($(am__cd) $(top_srcdir) && $(AUTOHEADER))
rm -f stamp-h1
touch $@
distclean-hdr:
-rm -f target-config.h stamp-h1
install-toolexeclibLTLIBRARIES: $(toolexeclib_LTLIBRARIES)
@$(NORMAL_INSTALL)
@list='$(toolexeclib_LTLIBRARIES)'; test -n "$(toolexeclibdir)" || list=; \
list2=; for p in $$list; do \
if test -f $$p; then \
list2="$$list2 $$p"; \
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)"; \
}
uninstall-toolexeclibLTLIBRARIES:
@$(NORMAL_UNINSTALL)
@list='$(toolexeclib_LTLIBRARIES)'; test -n "$(toolexeclibdir)" || list=; \
for p in $$list; do \
$(am__strip_dir) \
echo " $(LIBTOOL) $(AM_LIBTOOLFLAGS) $(LIBTOOLFLAGS) --mode=uninstall rm -f '$(DESTDIR)$(toolexeclibdir)/$$f'"; \
$(LIBTOOL) $(AM_LIBTOOLFLAGS) $(LIBTOOLFLAGS) --mode=uninstall rm -f "$(DESTDIR)$(toolexeclibdir)/$$f"; \
done
clean-toolexeclibLTLIBRARIES:
-test -z "$(toolexeclib_LTLIBRARIES)" || rm -f $(toolexeclib_LTLIBRARIES)
@list='$(toolexeclib_LTLIBRARIES)'; \
locs=`for p in $$list; do echo $$p; done | \
sed 's|^[^/]*$$|.|; s|/[^/]*$$||; s|$$|/so_locations|' | \
sort -u`; \
test -z "$$locs" || { \
echo rm -f $${locs}; \
rm -f $${locs}; \
}
rt/$(am__dirstamp):
@$(MKDIR_P) rt
@: > rt/$(am__dirstamp)
rt/$(DEPDIR)/$(am__dirstamp):
@$(MKDIR_P) rt/$(DEPDIR)
@: > rt/$(DEPDIR)/$(am__dirstamp)
rt/arithmetic.lo: rt/$(am__dirstamp) rt/$(DEPDIR)/$(am__dirstamp)
rt/atomics.lo: rt/$(am__dirstamp) rt/$(DEPDIR)/$(am__dirstamp)
rt/bitstring.lo: rt/$(am__dirstamp) rt/$(DEPDIR)/$(am__dirstamp)
rt/fbarrier.lo: rt/$(am__dirstamp) rt/$(DEPDIR)/$(am__dirstamp)
rt/fp16.lo: rt/$(am__dirstamp) rt/$(DEPDIR)/$(am__dirstamp)
rt/misc.lo: rt/$(am__dirstamp) rt/$(DEPDIR)/$(am__dirstamp)
rt/multimedia.lo: rt/$(am__dirstamp) rt/$(DEPDIR)/$(am__dirstamp)
rt/queue.lo: rt/$(am__dirstamp) rt/$(DEPDIR)/$(am__dirstamp)
rt/sat_arithmetic.lo: rt/$(am__dirstamp) rt/$(DEPDIR)/$(am__dirstamp)
rt/segment.lo: rt/$(am__dirstamp) rt/$(DEPDIR)/$(am__dirstamp)
rt/workitems.lo: rt/$(am__dirstamp) rt/$(DEPDIR)/$(am__dirstamp)
rt/fibers.lo: rt/$(am__dirstamp) rt/$(DEPDIR)/$(am__dirstamp)
libhsail-rt.la: $(libhsail_rt_la_OBJECTS) $(libhsail_rt_la_DEPENDENCIES) $(EXTRA_libhsail_rt_la_DEPENDENCIES)
$(AM_V_CCLD)$(libhsail_rt_la_LINK) -rpath $(toolexeclibdir) $(libhsail_rt_la_OBJECTS) $(libhsail_rt_la_LIBADD) $(LIBS)
mostlyclean-compile:
-rm -f *.$(OBJEXT)
-rm -f rt/*.$(OBJEXT)
-rm -f rt/*.lo
distclean-compile:
-rm -f *.tab.c
@AMDEP_TRUE@@am__include@ @am__quote@rt/$(DEPDIR)/arithmetic.Plo@am__quote@
@AMDEP_TRUE@@am__include@ @am__quote@rt/$(DEPDIR)/atomics.Plo@am__quote@
@AMDEP_TRUE@@am__include@ @am__quote@rt/$(DEPDIR)/bitstring.Plo@am__quote@
@AMDEP_TRUE@@am__include@ @am__quote@rt/$(DEPDIR)/fbarrier.Plo@am__quote@
@AMDEP_TRUE@@am__include@ @am__quote@rt/$(DEPDIR)/fibers.Plo@am__quote@
@AMDEP_TRUE@@am__include@ @am__quote@rt/$(DEPDIR)/fp16.Plo@am__quote@
@AMDEP_TRUE@@am__include@ @am__quote@rt/$(DEPDIR)/misc.Plo@am__quote@
@AMDEP_TRUE@@am__include@ @am__quote@rt/$(DEPDIR)/multimedia.Plo@am__quote@
@AMDEP_TRUE@@am__include@ @am__quote@rt/$(DEPDIR)/queue.Plo@am__quote@
@AMDEP_TRUE@@am__include@ @am__quote@rt/$(DEPDIR)/sat_arithmetic.Plo@am__quote@
@AMDEP_TRUE@@am__include@ @am__quote@rt/$(DEPDIR)/segment.Plo@am__quote@
@AMDEP_TRUE@@am__include@ @am__quote@rt/$(DEPDIR)/workitems.Plo@am__quote@
.c.o:
@am__fastdepCC_TRUE@ $(AM_V_CC)depbase=`echo $@ | sed 's|[^/]*$$|$(DEPDIR)/&|;s|\.o$$||'`;\
@am__fastdepCC_TRUE@ $(COMPILE) -MT $@ -MD -MP -MF $$depbase.Tpo -c -o $@ $< &&\
@am__fastdepCC_TRUE@ $(am__mv) $$depbase.Tpo $$depbase.Po
@AMDEP_TRUE@@am__fastdepCC_FALSE@ $(AM_V_CC)source='$<' object='$@' libtool=no @AMDEPBACKSLASH@
@AMDEP_TRUE@@am__fastdepCC_FALSE@ DEPDIR=$(DEPDIR) $(CCDEPMODE) $(depcomp) @AMDEPBACKSLASH@
@am__fastdepCC_FALSE@ $(AM_V_CC@am__nodep@)$(COMPILE) -c -o $@ $<
.c.obj:
@am__fastdepCC_TRUE@ $(AM_V_CC)depbase=`echo $@ | sed 's|[^/]*$$|$(DEPDIR)/&|;s|\.obj$$||'`;\
@am__fastdepCC_TRUE@ $(COMPILE) -MT $@ -MD -MP -MF $$depbase.Tpo -c -o $@ `$(CYGPATH_W) '$<'` &&\
@am__fastdepCC_TRUE@ $(am__mv) $$depbase.Tpo $$depbase.Po
@AMDEP_TRUE@@am__fastdepCC_FALSE@ $(AM_V_CC)source='$<' object='$@' libtool=no @AMDEPBACKSLASH@
@AMDEP_TRUE@@am__fastdepCC_FALSE@ DEPDIR=$(DEPDIR) $(CCDEPMODE) $(depcomp) @AMDEPBACKSLASH@
@am__fastdepCC_FALSE@ $(AM_V_CC@am__nodep@)$(COMPILE) -c -o $@ `$(CYGPATH_W) '$<'`
.c.lo:
@am__fastdepCC_TRUE@ $(AM_V_CC)depbase=`echo $@ | sed 's|[^/]*$$|$(DEPDIR)/&|;s|\.lo$$||'`;\
@am__fastdepCC_TRUE@ $(LTCOMPILE) -MT $@ -MD -MP -MF $$depbase.Tpo -c -o $@ $< &&\
@am__fastdepCC_TRUE@ $(am__mv) $$depbase.Tpo $$depbase.Plo
@AMDEP_TRUE@@am__fastdepCC_FALSE@ $(AM_V_CC)source='$<' object='$@' libtool=yes @AMDEPBACKSLASH@
@AMDEP_TRUE@@am__fastdepCC_FALSE@ DEPDIR=$(DEPDIR) $(CCDEPMODE) $(depcomp) @AMDEPBACKSLASH@
@am__fastdepCC_FALSE@ $(AM_V_CC@am__nodep@)$(LTCOMPILE) -c -o $@ $<
mostlyclean-libtool:
-rm -f *.lo
clean-libtool:
-rm -rf .libs _libs
-rm -rf rt/.libs rt/_libs
distclean-libtool:
-rm -f libtool config.lt
ID: $(am__tagged_files)
$(am__define_uniq_tagged_files); mkid -fID $$unique
tags: tags-am
TAGS: tags
tags-am: $(TAGS_DEPENDENCIES) $(am__tagged_files)
set x; \
here=`pwd`; \
$(am__define_uniq_tagged_files); \
shift; \
if test -z "$(ETAGS_ARGS)$$*$$unique"; then :; else \
test -n "$$unique" || unique=$$empty_fix; \
if test $$# -gt 0; then \
$(ETAGS) $(ETAGSFLAGS) $(AM_ETAGSFLAGS) $(ETAGS_ARGS) \
"$$@" $$unique; \
else \
$(ETAGS) $(ETAGSFLAGS) $(AM_ETAGSFLAGS) $(ETAGS_ARGS) \
$$unique; \
fi; \
fi
ctags: ctags-am
CTAGS: ctags
ctags-am: $(TAGS_DEPENDENCIES) $(am__tagged_files)
$(am__define_uniq_tagged_files); \
test -z "$(CTAGS_ARGS)$$unique" \
|| $(CTAGS) $(CTAGSFLAGS) $(AM_CTAGSFLAGS) $(CTAGS_ARGS) \
$$unique
GTAGS:
here=`$(am__cd) $(top_builddir) && pwd` \
&& $(am__cd) $(top_srcdir) \
&& gtags -i $(GTAGS_ARGS) "$$here"
cscope: cscope.files
test ! -s cscope.files \
|| $(CSCOPE) -b -q $(AM_CSCOPEFLAGS) $(CSCOPEFLAGS) -i cscope.files $(CSCOPE_ARGS)
clean-cscope:
-rm -f cscope.files
cscope.files: clean-cscope cscopelist
cscopelist: cscopelist-am
cscopelist-am: $(am__tagged_files)
list='$(am__tagged_files)'; \
case "$(srcdir)" in \
[\\/]* | ?:[\\/]*) sdir="$(srcdir)" ;; \
*) sdir=$(subdir)/$(srcdir) ;; \
esac; \
for i in $$list; do \
if test -f "$$i"; then \
echo "$(subdir)/$$i"; \
else \
echo "$$sdir/$$i"; \
fi; \
done >> $(top_builddir)/cscope.files
distclean-tags:
-rm -f TAGS ID GTAGS GRTAGS GSYMS GPATH tags
-rm -f cscope.out cscope.in.out cscope.po.out cscope.files
check-am: all-am
check: check-am
all-am: Makefile $(LTLIBRARIES) target-config.h
installdirs:
for dir in "$(DESTDIR)$(toolexeclibdir)"; do \
test -z "$$dir" || $(MKDIR_P) "$$dir"; \
done
install: install-am
install-exec: install-exec-am
install-data: install-data-am
uninstall: uninstall-am
install-am: all-am
@$(MAKE) $(AM_MAKEFLAGS) install-exec-am install-data-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
mostlyclean-generic:
clean-generic:
distclean-generic:
-test -z "$(CONFIG_CLEAN_FILES)" || rm -f $(CONFIG_CLEAN_FILES)
-test . = "$(srcdir)" || test -z "$(CONFIG_CLEAN_VPATH_FILES)" || rm -f $(CONFIG_CLEAN_VPATH_FILES)
-rm -f rt/$(DEPDIR)/$(am__dirstamp)
-rm -f rt/$(am__dirstamp)
maintainer-clean-generic:
@echo "This command is intended for maintainers to use"
@echo "it deletes files that may require special tools to rebuild."
clean: clean-am
clean-am: clean-generic clean-libtool clean-toolexeclibLTLIBRARIES \
mostlyclean-am
distclean: distclean-am
-rm -f $(am__CONFIG_DISTCLEAN_FILES)
-rm -rf rt/$(DEPDIR)
-rm -f Makefile
distclean-am: clean-am distclean-compile distclean-generic \
distclean-hdr distclean-libtool distclean-tags
dvi: dvi-am
dvi-am:
html: html-am
html-am:
info: info-am
info-am:
install-data-am:
install-dvi: install-dvi-am
install-dvi-am:
install-exec-am: install-toolexeclibLTLIBRARIES
install-html: install-html-am
install-html-am:
install-info: install-info-am
install-info-am:
install-man:
install-pdf: install-pdf-am
install-pdf-am:
install-ps: install-ps-am
install-ps-am:
installcheck-am:
maintainer-clean: maintainer-clean-am
-rm -f $(am__CONFIG_DISTCLEAN_FILES)
-rm -rf $(top_srcdir)/autom4te.cache
-rm -rf rt/$(DEPDIR)
-rm -f Makefile
maintainer-clean-am: distclean-am maintainer-clean-generic
mostlyclean: mostlyclean-am
mostlyclean-am: mostlyclean-compile mostlyclean-generic \
mostlyclean-libtool
pdf: pdf-am
pdf-am:
ps: ps-am
ps-am:
uninstall-am: uninstall-toolexeclibLTLIBRARIES
.MAKE: all install-am install-strip
.PHONY: CTAGS GTAGS TAGS all all-am am--refresh check check-am clean \
clean-cscope clean-generic clean-libtool \
clean-toolexeclibLTLIBRARIES cscope cscopelist-am ctags \
ctags-am distclean distclean-compile distclean-generic \
distclean-hdr distclean-libtool distclean-tags dvi dvi-am html \
html-am info info-am install install-am install-data \
install-data-am install-dvi install-dvi-am install-exec \
install-exec-am install-html install-html-am install-info \
install-info-am install-man install-pdf install-pdf-am \
install-ps install-ps-am install-strip \
install-toolexeclibLTLIBRARIES installcheck installcheck-am \
installdirs maintainer-clean maintainer-clean-generic \
mostlyclean mostlyclean-compile mostlyclean-generic \
mostlyclean-libtool pdf pdf-am ps ps-am tags tags-am uninstall \
uninstall-am uninstall-toolexeclibLTLIBRARIES
.PRECIOUS: Makefile
# Tell versions [3.59,3.63) of GNU make to not export all variables.
# Otherwise a system limit (for SysV at least) may be exceeded.
.NOEXPORT:

10
libhsail-rt/README
View File

@ -1,10 +0,0 @@
This library implements the agent-side runtime functionality required
to run HSA finalized programs produced by the BRIG frontend.
The library contains both the code required to run kernels on the agent
and also functions implementing more complex HSAIL instructions.
rt/workitems.c contains the runtime entry function that manages multiple
work-item execution using fibers or simple for-loops (in case of work groups
without barriers). Otherwise, the rest of the source files mostly contain
functions that typically map directly to HSAIL instructions.

1179
libhsail-rt/aclocal.m4 vendored
View File

File diff suppressed because it is too large Load Diff

17327
libhsail-rt/configure vendored
View File

File diff suppressed because it is too large Load Diff

160
libhsail-rt/configure.ac
View File

@ -1,160 +0,0 @@
# Starting point copied from libcilkrts:
#
# @copyright
# Copyright (C) 2011-2013, Intel Corporation
# All rights reserved.
#
# @copyright
# Redistribution and use in source and binary forms, with or without
# modification, are permitted provided that the following conditions
# are met:
#
# * Redistributions of source code must retain the above copyright
# notice, this list of conditions and the following disclaimer.
# * Redistributions in binary form must reproduce the above copyright
# notice, this list of conditions and the following disclaimer in
# the documentation and/or other materials provided with the
# distribution.
# * Neither the name of Intel Corporation nor the names of its
# contributors may be used to endorse or promote products derived
# from this software without specific prior written permission.
#
# @copyright
# THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS
# "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT
# LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR
# A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT
# HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT,
# INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
# BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS
# OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED
# AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT
# LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY
# WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE
# POSSIBILITY OF SUCH DAMAGE.
AC_INIT([phsa HSAIL runtime library], [1.0], [pekka.jaaskelainen@parmance.com])
# Needed to define ${target}. Needs to be very early to avoid annoying
# warning about calling AC_ARG_PROGRAM before AC_CANONICAL_SYSTEM
AC_CANONICAL_SYSTEM
target_alias=${target_alias-$host_alias}
AC_SUBST(target_alias)
AM_INIT_AUTOMAKE([1.11.6 foreign no-dist])
AM_MAINTAINER_MODE
AC_PROG_CC
AC_PROG_CXX
# AC_PROG_LIBTOOL
AC_CONFIG_FILES([Makefile])
if test "${multilib}" = "yes"; then
multilib_arg="--enable-multilib"
else
multilib_arg=
fi
AC_MSG_CHECKING([for --enable-version-specific-runtime-libs])
AC_ARG_ENABLE([version-specific-runtime-libs],
AC_HELP_STRING([--enable-version-specific-runtime-libs],
[Specify that runtime libraries should be installed in a compi
ler-specific directory]),
[case "$enableval" in
yes) enable_version_specific_runtime_libs=yes ;;
no) enable_version_specific_runtime_libs=no ;;
*) AC_MSG_ERROR([Unknown argument to enable/disable version-specific libs
]);;
esac],
[enable_version_specific_runtime_libs=no])
AC_MSG_RESULT($enable_version_specific_runtime_libs)
GCC_WITH_TOOLEXECLIBDIR
# Calculate toolexeclibdir
# Also toolexecdir, though it's only used in toolexeclibdir
case ${enable_version_specific_runtime_libs} in
yes)
# Need the gcc compiler version to know where to install libraries
# and header files if --enable-version-specific-runtime-libs option
# is selected.
toolexecdir='$(libdir)/gcc/$(target_alias)'
toolexeclibdir='$(toolexecdir)/$(gcc_version)$(MULTISUBDIR)'
;;
no)
if test -n "$with_cross_host" &&
test x"$with_cross_host" != x"no"; then
# Install a library built with a cross compiler in tooldir, not libdir.
toolexecdir='$(exec_prefix)/$(target_alias)'
case ${with_toolexeclibdir} in
no)
toolexeclibdir='$(toolexecdir)/lib'
;;
*)
toolexeclibdir=${with_toolexeclibdir}
;;
esac
else
toolexecdir='$(libdir)/gcc-lib/$(target_alias)'
toolexeclibdir='$(libdir)'
fi
multi_os_directory=`$CC -print-multi-os-directory`
case $multi_os_directory in
.) ;; # Avoid trailing /.
*) toolexeclibdir=$toolexeclibdir/$multi_os_directory ;;
esac
;;
esac
# Set config_dir based on the target. config_dir specifies where to get
# target-specific files. The generic implementation is incomplete, but
# contains information on what's needed
case "${target}" in
x86_64-*-*)
config_dir="x86"
;;
i?86-*-*)
config_dir="x86"
;;
*)
config_dir="generic"
;;
esac
AC_SUBST(config_dir)
# We have linker scripts for appropriate operating systems
linux_linker_script=no
case "${host}" in
*-*-linux*)
linux_linker_script=yes
;;
esac
AM_CONDITIONAL(LINUX_LINKER_SCRIPT, test "$linux_linker_script" = "yes")
mac_linker_script=no
case "${host}" in
*-*-apple*)
mac_linker_script=yes
;;
esac
AM_CONDITIONAL(MAC_LINKER_SCRIPT, test "$mac_linker_script" = "yes")
AC_LIBTOOL_DLOPEN
AM_PROG_LIBTOOL
AC_SUBST(toolexecdir)
AC_SUBST(toolexeclibdir)
AC_CONFIG_HEADER(target-config.h)
AC_CHECK_SIZEOF([int])
AC_CHECK_SIZEOF([void*])
# Determine what GCC version number to use in filesystem paths.
GCC_BASE_VER
# Must be last
AC_OUTPUT

36
libhsail-rt/configure.tgt
View File

@ -1,36 +0,0 @@
# -*- shell-script -*-
# Copyright (C) 2012-2021 Free Software Foundation, Inc.
# Contributed by Pekka Jaaskelainen <pekka.jaaskelainen@parmance.com>
# for General Processor Tech.
#
# This file is part of the libhsail-rt.
#
# Permission is hereby granted, free of charge, to any person obtaining a
# copy of this software and associated documentation files
# (the "Software"), to deal in the Software without restriction, including
# without limitation the rights to use, copy, modify, merge, publish,
# distribute, sublicense, and/or sell copies of the Software, and to
# permit persons to whom the Software is furnished to do so, subject to
# the following conditions:
#
# The above copyright notice and this permission notice shall be included
# in all copies or substantial portions of the Software.
#
# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS
# OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
# MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.
# IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM,
# DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR
# OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE
# USE OR OTHER DEALINGS IN THE SOFTWARE.
# Disable the BRIG frontend and libhsail-rt on untested or known
# broken systems. Currently it has been tested only on x86_64 Linux
# of the upstream gcc targets. More targets shall be added after testing.
case "${target}" in
x86_64-*-linux* | i?86-*-linux*)
;;
*)
UNSUPPORTED=1
;;
esac

99
libhsail-rt/include/internal/fibers.h
View File

@ -1,99 +0,0 @@
/* fibers.h -- an extremely simple lightweight thread (fiber) implementation
Copyright (C) 2015-2021 Free Software Foundation, Inc.
Contributed by Pekka Jaaskelainen <pekka.jaaskelainen@parmance.com>
for General Processor Tech.
Permission is hereby granted, free of charge, to any person obtaining a
copy of this software and associated documentation files
(the "Software"), to deal in the Software without restriction, including
without limitation the rights to use, copy, modify, merge, publish,
distribute, sublicense, and/or sell copies of the Software, and to
permit persons to whom the Software is furnished to do so, subject to
the following conditions:
The above copyright notice and this permission notice shall be included
in all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS
OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.
IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM,
DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR
OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE
USE OR OTHER DEALINGS IN THE SOFTWARE.
*/
#ifndef PHSA_RT_FIBERS_H
#define PHSA_RT_FIBERS_H
#include <ucontext.h>
typedef enum
{
/* Ready to run. */
FIBER_STATUS_READY,
/* Exited by calling fiber_thread_exit. */
FIBER_STATUS_EXITED,
/* Joined by the main thread. */
FIBER_STATUS_JOINED
} fiber_status_t;
/* A light weight thread (fiber). */
struct fiber_s
{
ucontext_t context;
volatile fiber_status_t status;
struct fiber_s *next;
struct fiber_s *prev;
};
typedef struct fiber_s fiber_t;
typedef void (*fiber_function_t)(int, int);
/* Initializes the fiber with the start function given as the first
argument, and the argument to pass to the start function,
as the second. The allocated stack size is given as the last argument. */
void
fiber_init (fiber_t *fiber, fiber_function_t start_function, void *arg,
size_t stack_size, size_t stack_align);
/* Terminates the fiber execution from within the fiber itself. */
void
fiber_exit ();
/* Blocks until the given fiber returns. Frees the resources allocated
for the fiber. After join returns, the fiber itself can be deleted. */
void
fiber_join (fiber_t *fiber);
/* Co-operatively transfer execution turn to other fibers. */
void
fiber_yield ();
/* A multi-entry barrier. After the last fiber has reached the
barrier, it is automatically re-initialized to the threshold. */
typedef struct
{
/* The barrier participant count. */
volatile size_t threshold;
/* Number of fibers that have reached the barrier. */
volatile size_t reached;
/* Number of fibers that are waiting at the barrier. */
volatile size_t waiting_count;
} fiber_barrier_t;
/* Reach the given barrier. Blocks (co-operatively switches the execution
fibers) until all other parties have reached it. Returns 0 only in case
the calling fiber was the first one to return from the barrier. */
size_t
fiber_barrier_reach (fiber_barrier_t *barrier);
/* Initializes the given barrier. */
void
fiber_barrier_init (fiber_barrier_t *barrier, size_t threshold);
void *
fiber_int_args_to_ptr (int arg0, int arg1);
#endif

60
libhsail-rt/include/internal/phsa-queue-interface.h
View File

@ -1,60 +0,0 @@
/* phsa_queue_interface.h -- Definition for a minimalistic generic in-memory
representation of a user mode queue to be used with the phsa/gccbrig
implementation.
Copyright (C) 2015-2021 Free Software Foundation, Inc.
Contributed by Pekka Jaaskelainen <pekka.jaaskelainen@parmance.com>
for General Processor Tech.
Permission is hereby granted, free of charge, to any person obtaining a
copy of this software and associated documentation files
(the "Software"), to deal in the Software without restriction, including
without limitation the rights to use, copy, modify, merge, publish,
distribute, sublicense, and/or sell copies of the Software, and to
permit persons to whom the Software is furnished to do so, subject to
the following conditions:
The above copyright notice and this permission notice shall be included
in all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS
OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.
IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM,
DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR
OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE
USE OR OTHER DEALINGS IN THE SOFTWARE.
*/
#ifndef PHSA_QUEUE_INTERFACE_H
#define PHSA_QUEUE_INTERFACE_H
#ifdef __cplusplus
extern "C" {
#endif
#include <stdbool.h>
#include <stddef.h>
#include <stdint.h>
#include "hsa.h"
typedef __attribute__ ((aligned (64))) struct phsa_queue_s
{
/* An HSA Architectured Queue object. Must be in the beginning
of the struct to enable direct pointer casting between hsa_queue_
and phsa_queue_t. */
hsa_queue_t hsa_queue;
volatile uint64_t write_index;
volatile uint64_t read_index;
/* True if global mem addresses are 64b. */
uint64_t is_ptr64 : 1;
} phsa_queue_t;
#ifdef __cplusplus
}
#endif
#endif

94
libhsail-rt/include/internal/phsa-rt.h
View File

@ -1,94 +0,0 @@
/* phsa-rt.h -- Data structures and functions of the PHSA device side runtime
scheduler, and HSAIL built-ins.
Copyright (C) 2015-2021 Free Software Foundation, Inc.
Contributed by Pekka Jaaskelainen <pekka.jaaskelainen@parmance.com>
for General Processor Tech.
Permission is hereby granted, free of charge, to any person obtaining a
copy of this software and associated documentation files
(the "Software"), to deal in the Software without restriction, including
without limitation the rights to use, copy, modify, merge, publish,
distribute, sublicense, and/or sell copies of the Software, and to
permit persons to whom the Software is furnished to do so, subject to
the following conditions:
The above copyright notice and this permission notice shall be included
in all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS
OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.
IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM,
DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR
OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE
USE OR OTHER DEALINGS IN THE SOFTWARE.
*/
#ifndef PHSA_RT_H
#define PHSA_RT_H
#include <stdbool.h>
#include <stdint.h>
#include "hsa.h"
#define PHSA_MAX_WG_SIZE 1024 * 10
/* Pointer type for the public facing kernel launcher function generated
by gccbrig. This launches the actual kernel for all work groups and
work items in the grid. */
typedef void (*gccbrigKernelLauncherFunc) (void *context, void *);
/* Pointer type for kernel functions produced by gccbrig from the HSAIL.
This is private from outside the device binary and only called by
the launcher. */
typedef void (*gccbrigKernelFunc) (unsigned char *, void *, void *, uint32_t,
void *);
/* Context data that is passed to the kernel function, initialized
by the runtime to the current launch information. The data is
used by different id functions etc.
The struct is used by both the launcher and the targeted device,
thus the fields must have the same alignment/padding in both sides.
*/
typedef struct
{
/* Data set by the HSA Runtime's kernel launcher. */
hsa_kernel_dispatch_packet_t *dp;
size_t packet_id;
/* Data set by the device-side launcher. */
gccbrigKernelFunc kernel;
/* The range of a work groups this dispatch should execute. */
size_t wg_min_x;
size_t wg_min_y;
size_t wg_min_z;
size_t wg_max_x;
size_t wg_max_y;
size_t wg_max_z;
/* The barrier used to synch the work-items before executing a new WG. */
void *wg_start_barrier;
/* The barrier to wait at after executing a work-group. */
void *wg_completion_barrier;
/* The barrier used to synchronize WIs in case of the 'barrier' HSAIL
instruction. */
void *wg_sync_barrier;
/* This should be set to the flat address of the beginning of the group
segment. */
size_t group_segment_start_addr;
/* This must be set to the correct aligned flat address space location from
where the kernel can actually read its arguments. Might point to the
original global kernarg space. */
void *kernarg_addr;
} PHSAKernelLaunchData;
#endif

142
libhsail-rt/include/internal/workitems.h
View File

@ -1,142 +0,0 @@
/* workitems.h -- Types for context data passed as hidden parameters to special
built-ins.
Copyright (C) 2015-2021 Free Software Foundation, Inc.
Contributed by Pekka Jaaskelainen <pekka.jaaskelainen@parmance.com>
for General Processor Tech.
Permission is hereby granted, free of charge, to any person obtaining a
copy of this software and associated documentation files
(the "Software"), to deal in the Software without restriction, including
without limitation the rights to use, copy, modify, merge, publish,
distribute, sublicense, and/or sell copies of the Software, and to
permit persons to whom the Software is furnished to do so, subject to
the following conditions:
The above copyright notice and this permission notice shall be included
in all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS
OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.
IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM,
DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR
OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE
USE OR OTHER DEALINGS IN THE SOFTWARE.
*/
#ifndef PHSA_RT_WORKITEMS_H
#define PHSA_RT_WORKITEMS_H
/* As the simple fibers implementation relies only on ucontext, we can
assume is found by default as it is part of glibc. However, for partial
HSAIL support on platforms without having it available, the following define
can be undefined. */
#define HAVE_FIBERS
#ifdef HAVE_FIBERS
#include "fibers.h"
#endif
#include <stdint.h>
#include "phsa-rt.h"
/* Data identifying a single work-group instance. */
typedef struct
{
/* This is 1 in case there are more work groups to execute.
If 0, the work-item threads should finish themselves. */
int more_wgs;
/* If the local size does not evenly divide the grid size, will have
leftover WIs in the last execution. */
int leftover_wg;
int last_wg;
/* (Flat) pointer to the beginning of the group segment allocated
to the work-group. */
void *group_base_ptr;
/* The offset in the group memory for the kernel local group variables.
To support module scope group variables, there might be need to preseve
room for them in the beginning of the group segment. */
uint32_t initial_group_offset;
/* Similarly to the private segment that gets space allocated for all
WIs in the work-group. */
void *private_base_ptr;
uint32_t private_segment_total_size;
/* The first flat address of the group segment allocated for
the given work group. */
uint64_t group_segment_base_addr;
/* Offset from the beginning of the private segment to the start of
the previously allocated chunk of dynamic work-item memory (alloca)
by any WI in the WG.
Initially set to private_segment_total_size to denote no dynamic
allocations have been made. The dynamic allocations are done downwards
from the private segment end. */
uint32_t alloca_stack_p;
/* The position of the first word in the current function's alloca
stack frame. Initialized to point outside the private segment. */
uint32_t alloca_frame_p;
/* The group id of the currently executed WG. This is for fiber based
execution. The group ids are duplicated also to the per WI context
struct for simplified single pointer access in the GCCBRIG produced
code.
*/
uint32_t x;
uint32_t y;
uint32_t z;
} PHSAWorkGroup;
/* Data identifying a single work-item, passed to the work-item thread in case
of a fiber based work-group execution. */
typedef struct
{
/* NOTE: These members STARTing here should not be moved as they are
accessed directly by code emitted by BRIG FE. */
/* The local id of the current WI. */
uint32_t x;
uint32_t y;
uint32_t z;
/* The group id of the currently executed WG. */
uint32_t group_x;
uint32_t group_y;
uint32_t group_z;
/* The local size of a complete WG. */
uint32_t wg_size_x;
uint32_t wg_size_y;
uint32_t wg_size_z;
/* The local size of the current WG. */
uint32_t cur_wg_size_x;
uint32_t cur_wg_size_y;
uint32_t cur_wg_size_z;
/* NOTE: Fixed members END here. */
PHSAKernelLaunchData *launch_data;
/* Identifies and keeps book of the currently executed WG of the WI swarm. */
volatile PHSAWorkGroup *wg;
#ifdef HAVE_FIBERS
fiber_t fiber;
#endif
} __attribute__((packed)) PHSAWorkItem;
#endif

475
libhsail-rt/rt/arithmetic.c
View File

@ -1,475 +0,0 @@
/* arithmetic.c -- Builtins for HSAIL arithmetic instructions for which
there is no feasible direct gcc GENERIC expression.
Copyright (C) 2015-2021 Free Software Foundation, Inc.
Contributed by Pekka Jaaskelainen <pekka.jaaskelainen@parmance.com>
for General Processor Tech.
Permission is hereby granted, free of charge, to any person obtaining a
copy of this software and associated documentation files
(the "Software"), to deal in the Software without restriction, including
without limitation the rights to use, copy, modify, merge, publish,
distribute, sublicense, and/or sell copies of the Software, and to
permit persons to whom the Software is furnished to do so, subject to
the following conditions:
The above copyright notice and this permission notice shall be included
in all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS
OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.
IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM,
DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR
OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE
USE OR OTHER DEALINGS IN THE SOFTWARE.
*/
#include <stdio.h>
#include <stdint.h>
#include <limits.h>
#include <math.h>
#include <float.h>
/* HSAIL defines INT_MIN % -1 to be 0 while with C it's undefined,
and causes an overflow exception at least with gcc and C on IA-32. */
int32_t
__hsail_rem_s32 (int32_t dividend, int32_t divisor)
{
if (dividend == INT_MIN && divisor == -1)
return 0;
else
return dividend % divisor;
}
int64_t
__hsail_rem_s64 (int64_t dividend, int64_t divisor)
{
if (dividend == INT64_MIN && divisor == -1)
return 0;
else
return dividend % divisor;
}
/* HSAIL has defined behavior for min and max when one of the operands is
NaN: in that case the other operand is returned. In C and with gcc's
MIN_EXPR/MAX_EXPR, the returned operand is undefined. */
float
__hsail_min_f32 (float a, float b)
{
if (isnan (a))
return b;
else if (isnan (b))
return a;
else if (a == 0.0f && b == 0.0f)
return signbit (a) ? a : b;
else if (a > b)
return b;
else
return a;
}
double
__hsail_min_f64 (double a, double b)
{
if (isnan (a))
return b;
else if (isnan (b))
return a;
else if (a > b)
return b;
else
return a;
}
float
__hsail_max_f32 (float a, float b)
{
if (isnan (a))
return b;
else if (isnan (b))
return a;
else if (a == 0.0f && b == 0.0f && signbit (a))
return b;
else if (a < b)
return b;
else
return a;
}
double
__hsail_max_f64 (double a, double b)
{
if (isnan (a))
return b;
else if (isnan (b))
return a;
else if (a == 0.0 && b == 0.0 && signbit (a))
return b;
else if (a < b)
return b;
else
return a;
}
uint8_t
__hsail_cvt_zeroi_sat_u8_f32 (float a)
{
if (isnan (a))
return 0;
if (a >= (float) UINT8_MAX)
return UINT8_MAX;
else if (a <= 0.0f)
return 0;
return (uint8_t) a;
}
int8_t
__hsail_cvt_zeroi_sat_s8_f32 (float a)
{
if (isnan (a))
return 0;
if (a >= (float) INT8_MAX)
return INT8_MAX;
if (a <= (float) INT8_MIN)
return INT8_MIN;
return (int8_t) a;
}
uint16_t
__hsail_cvt_zeroi_sat_u16_f32 (float a)
{
if (isnan (a))
return 0;
if (a >= (float) UINT16_MAX)
return UINT16_MAX;
else if (a <= 0.0f)
return 0;
return (uint16_t) a;
}
int16_t
__hsail_cvt_zeroi_sat_s16_f32 (float a)
{
if (isnan (a))
return 0;
if (a >= (float) INT16_MAX)
return INT16_MAX;
if (a <= (float) INT16_MIN)
return INT16_MIN;
return (int16_t) a;
}
uint32_t
__hsail_cvt_zeroi_sat_u32_f32 (float a)
{
if (isnan (a))
return 0;
if (a >= (float) UINT32_MAX)
return UINT32_MAX;
else if (a <= 0.0f)
return 0;
return (uint32_t) a;
}
int32_t
__hsail_cvt_zeroi_sat_s32_f32 (float a)
{
if (isnan (a))
return 0;
if (a >= (float) INT32_MAX)
return INT32_MAX;
if (a <= (float) INT32_MIN)
return INT32_MIN;
return (int32_t) a;
}
uint64_t
__hsail_cvt_zeroi_sat_u64_f32 (float a)
{
if (isnan (a))
return 0;
if (a >= (float) UINT64_MAX)
return UINT64_MAX;
else if (a <= 0.0f)
return 0;
return (uint64_t) a;
}
int64_t
__hsail_cvt_zeroi_sat_s64_f32 (float a)
{
if (isnan (a))
return 0;
if (a >= (float) INT64_MAX)
return INT64_MAX;
if (a <= (float) INT64_MIN)
return INT64_MIN;
return (int64_t) a;
}
uint8_t
__hsail_cvt_zeroi_sat_u8_f64 (double a)
{
if (isnan (a))
return 0;
if (a >= (double) UINT8_MAX)
return UINT8_MAX;
else if (a <= 0.0f)
return 0;
return (uint8_t) a;
}
int8_t
__hsail_cvt_zeroi_sat_s8_f64 (double a)
{
if (isnan (a))
return 0;
if (a >= (double) INT8_MAX)
return INT8_MAX;
if (a <= (double) INT8_MIN)
return INT8_MIN;
return (int8_t) a;
}
uint16_t
__hsail_cvt_zeroi_sat_u16_f64 (double a)
{
if (isnan (a))
return 0;
if (a >= (double) UINT16_MAX)
return UINT16_MAX;
else if (a <= 0.0f)
return 0;
return (uint16_t) a;
}
int16_t
__hsail_cvt_zeroi_sat_s16_f64 (double a)
{
if (isnan (a))
return 0;
if (a >= (double) INT16_MAX)
return INT16_MAX;
if (a <= (double) INT16_MIN)
return INT16_MIN;
return (int16_t) a;
}
uint32_t
__hsail_cvt_zeroi_sat_u32_f64 (double a)
{
if (isnan (a))
return 0;
if (a >= (double) UINT32_MAX)
return UINT32_MAX;
else if (a <= 0.0f)
return 0;
return (uint32_t) a;
}
int32_t
__hsail_cvt_zeroi_sat_s32_f64 (double a)
{
if (isnan (a))
return 0;
if (a >= (double) INT32_MAX)
return INT32_MAX;
if (a <= (double) INT32_MIN)
return INT32_MIN;
return (int32_t) a;
}
uint64_t
__hsail_cvt_zeroi_sat_u64_f64 (double a)
{
if (isnan (a))
return 0;
if (a >= (double) UINT64_MAX)
return UINT64_MAX;
else if (a <= 0.0f)
return 0;
return (uint64_t) a;
}
int64_t
__hsail_cvt_zeroi_sat_s64_f64 (double a)
{
if (isnan (a))
return 0;
if (a >= (double) INT64_MAX)
return INT64_MAX;
if (a <= (double) INT64_MIN)
return INT64_MIN;
return (int64_t) a;
}
/* Flush the operand to zero in case it's a denormalized number.
Do not cause any exceptions in case of NaNs. */
float
__hsail_ftz_f32 (float a)
{
if (isnan (a) || isinf (a) || a == 0.0f)
return a;
if (a < 0.0f)
{
if (-a < FLT_MIN)
return -0.0f;
}
else
{
if (a < FLT_MIN)
return 0.0f;
}
return a;
}
#define F16_MIN (6.10e-5)
/* Flush the single precision operand to zero in case it's considered
a denormalized number in case it was a f16. Do not cause any exceptions
in case of NaNs. */
float
__hsail_ftz_f32_f16 (float a)
{
if (isnan (a) || isinf (a) || a == 0.0f)
return a;
if (a < 0.0f)
{
if (-a < F16_MIN)
return -0.0f;
}
else
{
if (a < F16_MIN)
return 0.0f;
}
return a;
}
double
__hsail_ftz_f64 (double a)
{
if (isnan (a) || isinf (a) || a == 0.0d)
return a;
if (a < 0.0d)
{
if (-a < DBL_MIN)
return -0.0d;
}
else
{
if (a < DBL_MIN)
return 0.0d;
}
return a;
}
uint32_t
__hsail_borrow_u32 (uint32_t a, uint32_t b)
{
return __builtin_sub_overflow_p (a, b, a);
}
uint64_t
__hsail_borrow_u64 (uint64_t a, uint64_t b)
{
return __builtin_sub_overflow_p (a, b, a);
}
uint32_t
__hsail_carry_u32 (uint32_t a, uint32_t b)
{
return __builtin_add_overflow_p (a, b, a);
}
uint64_t
__hsail_carry_u64 (uint64_t a, uint64_t b)
{
return __builtin_add_overflow_p (a, b, a);
}
float
__hsail_fract_f32 (float a)
{
int exp;
if (isinf (a))
return signbit (a) == 0 ? 0.0f : -0.0f;
if (isnan (a) || a == 0.0f)
return a;
else
return fminf (a - floorf (a), 0x1.fffffep-1f);
}
double
__hsail_fract_f64 (double a)
{
int exp;
if (isinf (a))
return 0.0f * isinf (a);
if (isnan (a) || a == 0.0f)
return a;
else
return fmin (a - floor (a), 0x1.fffffffffffffp-1d);
}
uint32_t
__hsail_class_f32 (float a, uint32_t flags)
{
return (flags & 0x0001 && isnan (a) && !(*(uint32_t *) &a & (1ul << 22)))
|| (flags & 0x0002 && isnan (a) && (*(uint32_t *) &a & (1ul << 22)))
|| (flags & 0x0004 && isinf (a) && a < 0.0f)
|| (flags & 0x0008 && isnormal (a) && signbit (a))
|| (flags & 0x0010 && a < 0.0f && a > -FLT_MIN)
|| (flags & 0x0020 && a == 0.0f && signbit (a))
|| (flags & 0x0040 && a == 0.0f && !signbit (a))
|| (flags & 0x0080 && a > 0.0f && a < FLT_MIN)
|| (flags & 0x0100 && isnormal (a) && !signbit (a))
|| (flags & 0x0200 && isinf (a) && a >= 0.0f);
}
uint32_t
__hsail_class_f64 (double a, uint32_t flags)
{
return (flags & 0x0001 && isnan (a) && !(*(uint64_t *) &a & (1ul << 51)))
|| (flags & 0x0002 && isnan (a) && (*(uint64_t *) &a & (1ul << 51)))
|| (flags & 0x0004 && isinf (a) && a < 0.0f)
|| (flags & 0x0008 && isnormal (a) && signbit (a))
|| (flags & 0x0010 && a < 0.0f && a > -FLT_MIN)
|| (flags & 0x0020 && a == 0.0f && signbit (a))
|| (flags & 0x0040 && a == 0.0f && !signbit (a))
|| (flags & 0x0080 && a > 0.0f && a < FLT_MIN)
|| (flags & 0x0100 && isnormal (a) && !signbit (a))
|| (flags & 0x0200 && isinf (a) && a >= 0.0f);
}
/* 'class' for a f32-converted f16 which should otherwise be treated like f32
except for its limits. */
uint32_t
__hsail_class_f32_f16 (float a, uint32_t flags)
{
return (flags & 0x0001 && isnan (a) && !(*(uint32_t *) &a & 0x40000000))
|| (flags & 0x0002 && isnan (a) && (*(uint32_t *) &a & 0x40000000))
|| (flags & 0x0004 && isinf (a) && a < 0.0f)
|| (flags & 0x0008 && a != 0.0f && !isinf (a) && !isnan (a)
&& a <= -F16_MIN)
|| (flags & 0x0010 && a != 0.0f && !isinf (a) && !isnan (a) && a < 0.0f
&& a > -F16_MIN)
|| (flags & 0x0020 && a == 0.0f && signbit (a))
|| (flags & 0x0040 && a == 0.0f && !signbit (a))
|| (flags & 0x0080 && a != 0.0f && !isinf (a) && !isnan (a) && a > 0.0f
&& a < F16_MIN)
|| (flags & 0x0100 && a != 0.0f && !isinf (a) && !isnan (a)
&& a >= F16_MIN)
|| (flags & 0x0200 && isinf (a) && a >= 0.0f);
}

115
libhsail-rt/rt/atomics.c
View File

@ -1,115 +0,0 @@
/* atomic.c -- Builtins for HSAIL atomic instructions for which
there is no feasible direct gcc GENERIC expression.
Copyright (C) 2015-2021 Free Software Foundation, Inc.
Contributed by Pekka Jaaskelainen <pekka.jaaskelainen@parmance.com>
for General Processor Tech.
Permission is hereby granted, free of charge, to any person obtaining a
copy of this software and associated documentation files
(the "Software"), to deal in the Software without restriction, including
without limitation the rights to use, copy, modify, merge, publish,
distribute, sublicense, and/or sell copies of the Software, and to
permit persons to whom the Software is furnished to do so, subject to
the following conditions:
The above copyright notice and this permission notice shall be included
in all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS
OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.
IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM,
DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR
OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE
USE OR OTHER DEALINGS IN THE SOFTWARE.
*/
#include <stdint.h>
#include <stdio.h>
#define DO_ATOMICALLY(T, OPERATION) \
int done = 0; \
T old_value; \
T new_value; \
while (!done) \
{ \
old_value = *ptr; \
new_value = OPERATION; \
done = __sync_bool_compare_and_swap (ptr, old_value, new_value); \
} \
return old_value
int32_t
__hsail_atomic_min_s32 (int32_t *ptr, int32_t a)
{
DO_ATOMICALLY (int32_t, (old_value < a) ? old_value : a);
}
int64_t
__hsail_atomic_min_s64 (int64_t *ptr, int64_t a)
{
DO_ATOMICALLY (int64_t, (old_value < a) ? old_value : a);
}
uint32_t
__hsail_atomic_min_u32 (uint32_t *ptr, uint32_t a)
{
DO_ATOMICALLY (uint32_t, (old_value < a) ? old_value : a);
}
uint64_t
__hsail_atomic_min_u64 (uint64_t *ptr, uint64_t a)
{
DO_ATOMICALLY (uint64_t, (old_value < a) ? old_value : a);
}
uint32_t
__hsail_atomic_max_u32 (uint32_t *ptr, uint32_t a)
{
DO_ATOMICALLY (uint32_t, (old_value > a) ? old_value : a);
}
int32_t
__hsail_atomic_max_s32 (int32_t *ptr, int32_t a)
{
DO_ATOMICALLY (int32_t, (old_value > a) ? old_value : a);
}
uint64_t
__hsail_atomic_max_u64 (uint64_t *ptr, uint64_t a)
{
DO_ATOMICALLY (uint64_t, (old_value > a) ? old_value : a);
}
int64_t
__hsail_atomic_max_s64 (int64_t *ptr, int64_t a)
{
DO_ATOMICALLY (int64_t, (old_value > a) ? old_value : a);
}
uint32_t
__hsail_atomic_wrapinc_u32 (uint32_t *ptr, uint32_t a)
{
DO_ATOMICALLY (uint32_t, (old_value >= a) ? 0 : (old_value + 1));
}
uint64_t
__hsail_atomic_wrapinc_u64 (uint64_t *ptr, uint64_t a)
{
DO_ATOMICALLY (uint64_t, (old_value >= a) ? 0 : (old_value + 1));
}
uint32_t
__hsail_atomic_wrapdec_u32 (uint32_t *ptr, uint32_t a)
{
DO_ATOMICALLY (uint32_t,
((old_value == 0) || (old_value > a)) ? a : (old_value - 1));
}
uint64_t
__hsail_atomic_wrapdec_u64 (uint64_t *ptr, uint64_t a)
{
DO_ATOMICALLY (uint64_t,
((old_value == 0) || (old_value > a)) ? a : (old_value - 1));
}

190
libhsail-rt/rt/bitstring.c
View File

@ -1,190 +0,0 @@
/* bitstring.c -- Builtins for HSAIL bitstring instructions.
Copyright (C) 2015-2021 Free Software Foundation, Inc.
Contributed by Pekka Jaaskelainen <pekka.jaaskelainen@parmance.com>
for General Processor Tech.
Permission is hereby granted, free of charge, to any person obtaining a
copy of this software and associated documentation files
(the "Software"), to deal in the Software without restriction, including
without limitation the rights to use, copy, modify, merge, publish,
distribute, sublicense, and/or sell copies of the Software, and to
permit persons to whom the Software is furnished to do so, subject to
the following conditions:
The above copyright notice and this permission notice shall be included
in all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS
OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.
IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM,
DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR
OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE
USE OR OTHER DEALINGS IN THE SOFTWARE.
*/
#include <stdint.h>
#include <limits.h>
#define BITEXTRACT(DEST_TYPE, SRC0, SRC1, SRC2) \
uint32_t offset = SRC1 & (sizeof (DEST_TYPE) * 8 - 1); \
uint32_t width = SRC2 & (sizeof (DEST_TYPE) * 8 - 1); \
if (width == 0) \
return 0; \
else \
return (SRC0 << (sizeof (DEST_TYPE) * 8 - width - offset)) \
>> (sizeof (DEST_TYPE) * 8 - width)
uint32_t
__hsail_bitextract_u32 (uint32_t src0, uint32_t src1, uint32_t src2)
{
BITEXTRACT (uint32_t, src0, src1, src2);
}
int32_t
__hsail_bitextract_s32 (int32_t src0, uint32_t src1, uint32_t src2)
{
BITEXTRACT (int32_t, src0, src1, src2);
}
uint64_t
__hsail_bitextract_u64 (uint64_t src0, uint32_t src1, uint32_t src2)
{
BITEXTRACT (uint64_t, src0, src1, src2);
}
int64_t
__hsail_bitextract_s64 (int64_t src0, uint32_t src1, uint32_t src2)
{
BITEXTRACT (int64_t, src0, src1, src2);
}
#define BITINSERT(DEST_TYPE, SRC0, SRC1, SRC2, SRC3) \
uint32_t offset = SRC2 & (sizeof (DEST_TYPE) * 8 - 1); \
uint32_t width = SRC3 & (sizeof (DEST_TYPE) * 8 - 1); \
DEST_TYPE mask = ((DEST_TYPE) 1 << width) - 1; \
return (SRC0 & ~(mask << offset)) | ((SRC1 & mask) << offset)
uint32_t
__hsail_bitinsert_u32 (uint32_t src0, uint32_t src1, uint32_t src2,
uint32_t src3)
{
BITINSERT (uint32_t, src0, src1, src2, src3);
}
int64_t
__hsail_bitinsert_u64 (uint64_t src0, uint64_t src1, uint32_t src2,
uint32_t src3)
{
BITINSERT (uint64_t, src0, src1, src2, src3);
}
#define BITMASK(DEST_TYPE, SRC0, SRC1) \
uint32_t offset = SRC0 & (sizeof (DEST_TYPE) * 8 - 1); \
uint32_t width = SRC1 & (sizeof (DEST_TYPE) * 8 - 1); \
DEST_TYPE mask = ((DEST_TYPE) 1 << width) - 1; \
return mask << offset
uint32_t
__hsail_bitmask_u32 (uint32_t src0, uint32_t src1)
{
BITMASK (uint32_t, src0, src1);
}
uint64_t
__hsail_bitmask_u64 (uint32_t src0, uint32_t src1)
{
BITMASK (uint64_t, src0, src1);
}
/* The dummy, but readable version from
http://graphics.stanford.edu/~seander/bithacks.html#BitReverseObvious
This (also) often maps to a single instruction in DSPs. */
#define BITREV(DEST_TYPE, SRC) \
DEST_TYPE v = SRC; \
DEST_TYPE r = v; \
int s = sizeof (SRC) * CHAR_BIT - 1; \
\
for (v >>= 1; v; v >>= 1) \
{ \
r <<= 1; \
r |= v & 1; \
s--; \
} \
return r << s
uint32_t
__hsail_bitrev_u32 (uint32_t src0)
{
BITREV (uint32_t, src0);
}
uint64_t
__hsail_bitrev_u64 (uint64_t src0)
{
BITREV (uint64_t, src0);
}
uint32_t
__hsail_bitselect_u32 (uint32_t src0, uint32_t src1, uint32_t src2)
{
return (src1 & src0) | (src2 & ~src0);
}
uint64_t
__hsail_bitselect_u64 (uint64_t src0, uint64_t src1, uint64_t src2)
{
return (src1 & src0) | (src2 & ~src0);
}
/* Due to the defined behavior with 0, we cannot use the gcc builtin
__builtin_clz* () directly. __builtin_ffs () has defined behavior, but
returns 0 while HSAIL requires to return -1. */
uint32_t
__hsail_firstbit_u32 (uint32_t src0)
{
if (src0 == 0)
return -1;
return __builtin_clz (src0);
}
uint32_t
__hsail_firstbit_s32 (int32_t src0)
{
uint32_t converted = src0 >= 0 ? src0 : ~src0;
return __hsail_firstbit_u32 (converted);
}
uint32_t
__hsail_firstbit_u64 (uint64_t src0)
{
if (src0 == 0)
return -1;
return __builtin_clzl (src0);
}
uint32_t
__hsail_firstbit_s64 (int64_t src0)
{
uint64_t converted = src0 >= 0 ? src0 : ~src0;
return __hsail_firstbit_u64 (converted);
}
uint32_t
__hsail_lastbit_u32 (uint32_t src0)
{
if (src0 == 0)
return -1;
return __builtin_ctz (src0);
}
uint32_t
__hsail_lastbit_u64 (uint64_t src0)
{
if (src0 == 0)
return -1;
return __builtin_ctzl (src0);
}

87
libhsail-rt/rt/fbarrier.c
View File

@ -1,87 +0,0 @@
/* fbarrier.c -- HSAIL fbarrier built-ins.
Copyright (C) 2015-2021 Free Software Foundation, Inc.
Contributed by Pekka Jaaskelainen <pekka.jaaskelainen@parmance.com>
for General Processor Tech.
Permission is hereby granted, free of charge, to any person obtaining a
copy of this software and associated documentation files
(the "Software"), to deal in the Software without restriction, including
without limitation the rights to use, copy, modify, merge, publish,
distribute, sublicense, and/or sell copies of the Software, and to
permit persons to whom the Software is furnished to do so, subject to
the following conditions:
The above copyright notice and this permission notice shall be included
in all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS
OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.
IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM,
DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR
OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE
USE OR OTHER DEALINGS IN THE SOFTWARE.
*/
#include <stdlib.h>
#include <signal.h>
#include "workitems.h"
#include "phsa-rt.h"
#ifdef HAVE_FIBERS
#include "fibers.h"
typedef fiber_barrier_t fbarrier;
void
__hsail_initfbar (uint32_t addr, PHSAWorkItem *wi)
{
fbarrier *fbar = (fbarrier *) (wi->wg->group_base_ptr + addr);
fbar->threshold = 0;
fbar->reached = 0;
fbar->waiting_count = 0;
}
void
__hsail_releasefbar (uint32_t addr, PHSAWorkItem *wi)
{
fbarrier *fbar = (fbarrier *) (wi->wg->group_base_ptr + addr);
fbar->threshold = 0;
fbar->reached = 0;
fbar->waiting_count = 0;
}
void
__hsail_joinfbar (uint32_t addr, PHSAWorkItem *wi)
{
fbarrier *fbar = (fbarrier *) (wi->wg->group_base_ptr + addr);
++fbar->threshold;
}
void
__hsail_leavefbar (uint32_t addr, PHSAWorkItem *wi)
{
fbarrier *fbar = (fbarrier *) (wi->wg->group_base_ptr + addr);
--fbar->threshold;
}
void
__hsail_waitfbar (uint32_t addr, PHSAWorkItem *wi)
{
fbarrier *fbar = (fbarrier *) (wi->wg->group_base_ptr + addr);
fiber_barrier_reach (fbar);
}
void
__hsail_arrivefbar (uint32_t addr, PHSAWorkItem *wi)
{
fbarrier *fbar = (fbarrier *) (wi->wg->group_base_ptr + addr);
++fbar->reached;
if (fbar->reached == fbar->threshold)
fbar->reached = 0;
}
#endif

220
libhsail-rt/rt/fibers.c
View File

@ -1,220 +0,0 @@
/* fibers.c -- extremely simple lightweight thread (fiber) implementation
Copyright (C) 2016-2021 Free Software Foundation, Inc.
Contributed by Pekka Jaaskelainen <pekka.jaaskelainen@parmance.com>
for General Processor Tech.
Copyright (C) 2015-2021 Free Software Foundation, Inc.
Contributed by Pekka Jaaskelainen <pekka.jaaskelainen@parmance.com>
for General Processor Tech.
Permission is hereby granted, free of charge, to any person obtaining a
copy of this software and associated documentation files
(the "Software"), to deal in the Software without restriction, including
without limitation the rights to use, copy, modify, merge, publish,
distribute, sublicense, and/or sell copies of the Software, and to
permit persons to whom the Software is furnished to do so, subject to
the following conditions:
The above copyright notice and this permission notice shall be included
in all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS
OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.
IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM,
DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR
OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE
USE OR OTHER DEALINGS IN THE SOFTWARE.
*/
#include <stdlib.h>
#include <stdio.h>
#include <stdint.h>
#include "target-config.h"
#include "fibers.h"
void
phsa_fatal_error (int code);
ucontext_t main_context;
/* The last fiber in the linked list. */
static fiber_t *tail_fiber = NULL;
/* The first fiber in the linked list. */
static fiber_t *head_fiber = NULL;
/* The fiber currently being executed. */
static fiber_t *current_fiber = NULL;
/* Makecontext accepts only integer arguments. We need to split the
pointer argument in case pointer does not fit into int. This helper
function can be used to restore the pointer from the arguments. */
void *
fiber_int_args_to_ptr (int arg0, int arg1)
{
void *ptr = NULL;
#if SIZEOF_VOIDP == 8 && SIZEOF_INT == 4
ptr = (void*)(((uint64_t) arg0 & (uint64_t) 0xFFFFFFFF)
| ((uint64_t) arg1 << 32));
#elif SIZEOF_VOIDP == 4 && SIZEOF_INT == 4
ptr = (void*)arg0;
#else
# error Unsupported pointer/int size.
#endif
return ptr;
}
void
fiber_init (fiber_t *fiber, fiber_function_t start_function, void *arg,
size_t stack_size, size_t stack_align)
{
int arg0, arg1;
if (getcontext (&fiber->context) != 0)
phsa_fatal_error (3);
if (posix_memalign (&fiber->context.uc_stack.ss_sp, stack_align, stack_size)
!= 0)
phsa_fatal_error (4);
fiber->context.uc_stack.ss_size = stack_size;
fiber->context.uc_link = &main_context;
/* makecontext () accepts only integer arguments. Split the
pointer argument to two args in the case pointer does not fit
into one int. */
#if SIZEOF_VOIDP == 8 && SIZEOF_INT == 4
arg0 = (int32_t) 0xFFFFFFFF & (uint64_t)arg;
arg1 = (int32_t) 0xFFFFFFFF & ((uint64_t)arg >> 32);
#elif SIZEOF_VOIDP == 4 && SIZEOF_INT == 4
arg0 = (int)arg;
arg1 = 0;
#else
# error Unsupported pointer/int size.
#endif
makecontext (&fiber->context, (void*)start_function, 2, arg0, arg1);
fiber->status = FIBER_STATUS_READY;
fiber->next = NULL;
fiber->prev = NULL;
/* Create a linked list of the created fibers. Append the new one at
the end. */
if (tail_fiber == NULL)
tail_fiber = fiber;
else
{
tail_fiber->next = fiber;
fiber->prev = tail_fiber;
tail_fiber = fiber;
}
if (head_fiber == NULL)
head_fiber = fiber;
}
void
fiber_exit ()
{
fiber_status_t old_status = current_fiber->status;
current_fiber->status = FIBER_STATUS_EXITED;
if (old_status == FIBER_STATUS_JOINED)
/* In case this thread has been joined, return back to the joiner. */
swapcontext (&current_fiber->context, &main_context);
else
/* In case the thread exited while being yielded from another thread,
switch back to another fiber. */
fiber_yield ();
}
void
fiber_join (fiber_t *fiber)
{
fiber_t *next_ready_fiber = NULL;
current_fiber = fiber;
if (fiber->status != FIBER_STATUS_EXITED)
{
fiber->status = FIBER_STATUS_JOINED;
while (fiber->status != FIBER_STATUS_EXITED)
swapcontext (&main_context, &fiber->context);
}
/* Remove the successfully joined fiber from the linked list so we won't
access it later (the fiber itself might be freed after the join). */
if (fiber->prev != NULL)
fiber->prev->next = fiber->next;
if (fiber->next != NULL)
fiber->next->prev = fiber->prev;
if (head_fiber == fiber)
head_fiber = fiber->next;
if (tail_fiber == fiber)
tail_fiber = fiber->prev;
free (fiber->context.uc_stack.ss_sp);
}
void
fiber_yield ()
{
fiber_t *next_ready_fiber = current_fiber;
if (current_fiber == head_fiber
&& current_fiber == tail_fiber)
{
/* If the last fiber exits independently, there is no
fiber to switch to. Switch to the main context in that
case. */
if (current_fiber->status == FIBER_STATUS_EXITED)
swapcontext (&current_fiber->context, &main_context);
}
do {
next_ready_fiber = next_ready_fiber->next != NULL
? next_ready_fiber->next : head_fiber;
} while (next_ready_fiber != current_fiber
&& next_ready_fiber->status == FIBER_STATUS_EXITED);
fiber_t *old_current_fiber = current_fiber;
current_fiber = next_ready_fiber;
swapcontext (&old_current_fiber->context, &next_ready_fiber->context);
}
size_t
fiber_barrier_reach (fiber_barrier_t *barrier)
{
/* Yield once to ensure that there are no fibers waiting for
a previous triggering of the barrier in the waiting_count
loop. This should release them before we update the reached
counter again. */
fiber_yield ();
barrier->reached++;
++barrier->waiting_count;
while (barrier->reached < barrier->threshold)
fiber_yield ();
--barrier->waiting_count;
/* Wait until all the fibers have reached this point. */
while (barrier->waiting_count > 0)
fiber_yield ();
/* Now all fibers have been released from the barrier waiting
loop. We can now safely reset the reach count for new triggering. */
if (barrier->reached > 0)
{
barrier->reached = 0;
return 0;
}
return 1;
}
void
fiber_barrier_init (fiber_barrier_t *barrier, size_t threshold)
{
barrier->threshold = threshold;
barrier->waiting_count = 0;
barrier->reached = 0;
}

135
libhsail-rt/rt/fp16.c
View File

@ -1,135 +0,0 @@
/* Half-float conversion routines. Code mostly borrowed from the ARM's
builtin function.
Copyright (C) 2008-2021 Free Software Foundation, Inc.
Contributed by CodeSourcery.
This file 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.
This file is distributed in the hope that it will be useful, but
WITHOUT ANY WARRANTY; without even the implied warranty of
MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU
General Public License for more details.
Under Section 7 of GPL version 3, you are granted additional
permissions described in the GCC Runtime Library Exception, version
3.1, as published by the Free Software Foundation.
You should have received a copy of the GNU General Public License and
a copy of the GCC Runtime Library Exception along with this program;
see the files COPYING3 and COPYING.RUNTIME respectively. If not, see
<http://www.gnu.org/licenses/>. */
static inline unsigned short
__gnu_f2h_internal (unsigned int a, int ieee)
{
unsigned short sign = (a >> 16) & 0x8000;
int aexp = (a >> 23) & 0xff;
unsigned int mantissa = a & 0x007fffff;
unsigned int mask;
unsigned int increment;
if (aexp == 0xff)
{
if (!ieee)
return sign;
if (mantissa == 0)
return sign | 0x7c00; /* Infinity. */
/* Remaining cases are NaNs. Convert SNaN to QNaN. */
return sign | 0x7e00 | (mantissa >> 13);
}
if (aexp == 0 && mantissa == 0)
return sign;
aexp -= 127;
/* Decimal point between bits 22 and 23. */
mantissa |= 0x00800000;
if (aexp < -14)
{
mask = 0x00ffffff;
if (aexp >= -25)
mask >>= 25 + aexp;
}
else
mask = 0x00001fff;
/* Round. */
if (mantissa & mask)
{
increment = (mask + 1) >> 1;
if ((mantissa & mask) == increment)
increment = mantissa & (increment << 1);
mantissa += increment;
if (mantissa >= 0x01000000)
{
mantissa >>= 1;
aexp++;
}
}
if (ieee)
{
if (aexp > 15)
return sign | 0x7c00;
}
else
{
if (aexp > 16)
return sign | 0x7fff;
}
if (aexp < -24)
return sign;
if (aexp < -14)
{
mantissa >>= -14 - aexp;
aexp = -14;
}
/* We leave the leading 1 in the mantissa, and subtract one
from the exponent bias to compensate. */
return sign | (((aexp + 14) << 10) + (mantissa >> 13));
}
static unsigned int
__gnu_h2f_internal (unsigned short a, int ieee)
{
unsigned int sign = (unsigned int) (a & 0x8000) << 16;
int aexp = (a >> 10) & 0x1f;
unsigned int mantissa = a & 0x3ff;
if (aexp == 0x1f && ieee)
return sign | 0x7f800000 | (mantissa << 13);
if (aexp == 0)
{
int shift;
if (mantissa == 0)
return sign;
shift = __builtin_clz (mantissa) - 21;
mantissa <<= shift;
aexp = -shift;
}
return sign | (((aexp + 0x70) << 23) + (mantissa << 13));
}
unsigned short
__hsail_f32_to_f16 (unsigned int a)
{
return __gnu_f2h_internal (a, 1);
}
unsigned int
__hsail_f16_to_f32 (unsigned short a)
{
return __gnu_h2f_internal (a, 1);
}

89
libhsail-rt/rt/misc.c
View File

@ -1,89 +0,0 @@
/* misc.c -- Builtins for HSAIL misc instructions.
Copyright (C) 2015-2021 Free Software Foundation, Inc.
Contributed by Pekka Jaaskelainen <pekka.jaaskelainen@parmance.com>
for General Processor Tech.
Permission is hereby granted, free of charge, to any person obtaining a
copy of this software and associated documentation files
(the "Software"), to deal in the Software without restriction, including
without limitation the rights to use, copy, modify, merge, publish,
distribute, sublicense, and/or sell copies of the Software, and to
permit persons to whom the Software is furnished to do so, subject to
the following conditions:
The above copyright notice and this permission notice shall be included
in all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS
OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.
IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM,
DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR
OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE
USE OR OTHER DEALINGS IN THE SOFTWARE.
*/
#include <stdint.h>
#include <time.h>
#include "workitems.h"
/* Return the monotonic clock as nanoseconds. */
uint64_t
__hsail_clock ()
{
struct timespec t;
clock_gettime (CLOCK_MONOTONIC, &t);
return (uint64_t) t.tv_sec * 1000000000 + (uint64_t) t.tv_nsec;
}
uint32_t
__hsail_cuid (PHSAWorkItem *wi)
{
/* All WIs are executed with a single compute unit (core/thread)
for now. */
return 0;
}
uint32_t
__hsail_maxcuid (PHSAWorkItem *wi)
{
/* All WIs are executed with a single compute unit (core/thread)
for now. */
return 0;
}
void
__hsail_debugtrap (uint32_t src, PHSAWorkItem *wi)
{
/* Could we produce a SIGTRAP signal here to drop to gdb
console, or similar? In any case, the execution of the
kernel should halt.
*/
return;
}
uint32_t
__hsail_groupbaseptr (PHSAWorkItem *wi)
{
return (uint32_t) (uintptr_t) (wi->wg->group_base_ptr
- wi->launch_data->group_segment_start_addr);
}
uint64_t
__hsail_kernargbaseptr_u64 (PHSAWorkItem *wi)
{
/* For now assume only a single kernarg allocation at a time.
Proper kernarg memory management to do. */
return (uint64_t) (uintptr_t) wi->launch_data->kernarg_addr;
}
uint32_t
__hsail_kernargbaseptr_u32 (PHSAWorkItem *wi)
{
/* For now assume only a single kernarg allocation at a time.
Proper kernarg memory management to do. */
return 0;
}

135
libhsail-rt/rt/multimedia.c
View File

@ -1,135 +0,0 @@
/* multimedia.c -- Builtins for HSAIL multimedia instructions.
Copyright (C) 2015-2021 Free Software Foundation, Inc.
Contributed by Pekka Jaaskelainen <pekka.jaaskelainen@parmance.com>
for General Processor Tech.
Permission is hereby granted, free of charge, to any person obtaining a
copy of this software and associated documentation files
(the "Software"), to deal in the Software without restriction, including
without limitation the rights to use, copy, modify, merge, publish,
distribute, sublicense, and/or sell copies of the Software, and to
permit persons to whom the Software is furnished to do so, subject to
the following conditions:
The above copyright notice and this permission notice shall be included
in all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS
OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.
IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM,
DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR
OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE
USE OR OTHER DEALINGS IN THE SOFTWARE.
*/
#include <math.h>
#include <stdint.h>
uint32_t
__hsail_bitalign (uint64_t lower, uint64_t upper, uint32_t shift_amount)
{
shift_amount = shift_amount & 31;
uint64_t packed_value = (upper << 32) | lower;
return (packed_value >> shift_amount) & 0xFFFFFFFF;
}
uint32_t
__hsail_bytealign (uint64_t lower, uint64_t upper, uint32_t shift_amount)
{
shift_amount = (shift_amount & 3) * 8;
uint64_t packed_value = (upper << 32) | lower;
return (packed_value >> shift_amount) & 0xFFFFFFFF;
}
uint32_t
__hsail_lerp (uint32_t a, uint32_t b, uint32_t c)
{
uint32_t e3
= (((((a >> 24) & 0xff) + ((b >> 24) & 0xff) + ((c >> 24) & 0x1)) / 2)
& 0xff)
<< 24;
uint32_t e2
= (((((a >> 16) & 0xff) + ((b >> 16) & 0xff) + ((c >> 16) & 0x1)) / 2)
& 0xff)
<< 16;
uint32_t e1
= (((((a >> 8) & 0xff) + ((b >> 8) & 0xff) + ((c >> 8) & 0x1)) / 2) & 0xff)
<< 8;
uint32_t e0 = (((a & 0xff) + (b & 0xff) + (c & 0x1)) / 2) & 0xff;
return e3 | e2 | e1 | e0;
}
static uint8_t
cvt_neari_sat_u8_f32 (float a)
{
if (isinf (a))
{
if (signbit (a)) return 0;
else return 255;
}
else if (isnan (a)) return 0;
else if (a < 0.0)
return 0;
else if (a > 255.0)
return 255;
else
return (uint8_t) a;
}
uint32_t
__hsail_packcvt (float a, float b, float c, float d)
{
return (uint32_t) cvt_neari_sat_u8_f32 (a)
| (uint32_t) cvt_neari_sat_u8_f32 (b) << 8
| (uint32_t) cvt_neari_sat_u8_f32 (c) << 16
| (uint32_t) cvt_neari_sat_u8_f32 (d) << 24;
}
float
__hsail_unpackcvt (uint32_t val, uint32_t index)
{
return (float) ((val >> (index * 8)) & 0xff);
}
static uint32_t
abs_diff (uint32_t a, uint32_t b)
{
if (a < b)
return b - a;
else
return a - b;
}
uint32_t
__hsail_sad_u8x4 (uint32_t a, uint32_t b, uint32_t add)
{
return abs_diff ((a >> 24) & 0xff, (b >> 24) & 0xff)
+ abs_diff ((a >> 16) & 0xff, (b >> 16) & 0xff)
+ abs_diff ((a >> 8) & 0xff, (b >> 8) & 0xff)
+ abs_diff ((a >> 0) & 0xff, (b >> 0) & 0xff) + add;
}
uint32_t
__hsail_sad_u16x2 (uint32_t a, uint32_t b, uint32_t add)
{
return abs_diff ((a >> 16) & 0xffff, (b >> 16) & 0xffff)
+ abs_diff ((a >> 0) & 0xffff, (b >> 0) & 0xffff) + add;
}
uint32_t
__hsail_sad_u32 (uint32_t a, uint32_t b, uint32_t add)
{
return abs_diff (a, b) + add;
}
uint32_t
__hsail_sadhi_u16x2_u8x4 (uint32_t a, uint32_t b, uint32_t add)
{
return (abs_diff ((a >> 24) & 0xff, (b >> 24) & 0xff) << 16)
+ (abs_diff ((a >> 16) & 0xff, (b >> 16) & 0xff) << 16)
+ (abs_diff ((a >> 8) & 0xff, (b >> 8) & 0xff) << 16)
+ (abs_diff ((a >> 0) & 0xff, (b >> 0) & 0xff) << 16) + add;
}

71
libhsail-rt/rt/queue.c
View File

@ -1,71 +0,0 @@
/* queue.c -- Builtins for HSAIL queue related instructions.
Copyright (C) 2015-2021 Free Software Foundation, Inc.
Contributed by Pekka Jaaskelainen <pekka.jaaskelainen@parmance.com>
for General Processor Tech.
Permission is hereby granted, free of charge, to any person obtaining a
copy of this software and associated documentation files
(the "Software"), to deal in the Software without restriction, including
without limitation the rights to use, copy, modify, merge, publish,
distribute, sublicense, and/or sell copies of the Software, and to
permit persons to whom the Software is furnished to do so, subject to
the following conditions:
The above copyright notice and this permission notice shall be included
in all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS
OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.
IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM,
DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR
OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE
USE OR OTHER DEALINGS IN THE SOFTWARE.
*/
#include "phsa-queue-interface.h"
uint64_t
__hsail_ldqueuereadindex (uint64_t queue_addr)
{
phsa_queue_t *queue = (phsa_queue_t *) (uintptr_t) queue_addr;
return queue->read_index;
}
uint64_t
__hsail_ldqueuewriteindex (uint64_t queue_addr)
{
phsa_queue_t *queue = (phsa_queue_t *) (uintptr_t) queue_addr;
return queue->write_index;
}
uint64_t
__hsail_addqueuewriteindex (uint64_t queue_addr, uint64_t value)
{
phsa_queue_t *queue = (phsa_queue_t *) (uintptr_t) queue_addr;
return __sync_fetch_and_add (&queue->write_index, value);
}
uint64_t
__hsail_casqueuewriteindex (uint64_t queue_addr, uint64_t cmp_value,
uint64_t new_value)
{
phsa_queue_t *queue = (phsa_queue_t *) (uintptr_t) queue_addr;
return __sync_val_compare_and_swap (&queue->write_index, cmp_value,
new_value);
}
void
__hsail_stqueuereadindex (uint64_t queue_addr, uint64_t value)
{
phsa_queue_t *queue = (phsa_queue_t *) (uintptr_t) queue_addr;
queue->read_index = value;
}
void
__hsail_stqueuewriteindex (uint64_t queue_addr, uint64_t value)
{
phsa_queue_t *queue = (phsa_queue_t *) (uintptr_t) queue_addr;
queue->write_index = value;
}

267
libhsail-rt/rt/sat_arithmetic.c
View File

@ -1,267 +0,0 @@
/* sat_arithmetic.c -- Builtins for HSAIL saturating arithmetic instructions.
Copyright (C) 2015-2021 Free Software Foundation, Inc.
Contributed by Pekka Jaaskelainen <pekka.jaaskelainen@parmance.com>
for General Processor Tech.
Permission is hereby granted, free of charge, to any person obtaining a
copy of this software and associated documentation files
(the "Software"), to deal in the Software without restriction, including
without limitation the rights to use, copy, modify, merge, publish,
distribute, sublicense, and/or sell copies of the Software, and to
permit persons to whom the Software is furnished to do so, subject to
the following conditions:
The above copyright notice and this permission notice shall be included
in all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS
OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.
IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM,
DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR
OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE
USE OR OTHER DEALINGS IN THE SOFTWARE.
*/
#include <stdint.h>
uint8_t
__hsail_sat_add_u8 (uint8_t a, uint8_t b)
{
uint16_t c = (uint16_t) a + (uint16_t) b;
if (c > UINT8_MAX)
return UINT8_MAX;
else
return c;
}
uint16_t
__hsail_sat_add_u16 (uint16_t a, uint16_t b)
{
uint32_t c = (uint32_t) a + (uint32_t) b;
if (c > UINT16_MAX)
return UINT16_MAX;
else
return c;
}
uint32_t
__hsail_sat_add_u32 (uint32_t a, uint32_t b)
{
uint32_t c;
if (__builtin_add_overflow (a, b, &c))
return UINT32_MAX;
return c;
}
uint64_t
__hsail_sat_add_u64 (uint64_t a, uint64_t b)
{
uint64_t c;
if (__builtin_add_overflow (a, b, &c))
return UINT64_MAX;
return c;
}
int8_t
__hsail_sat_add_s8 (int8_t a, int8_t b)
{
int16_t c = (int16_t) a + (int16_t) b;
if (c > INT8_MAX)
return INT8_MAX;
else if (c < INT8_MIN)
return INT8_MIN;
else
return c;
}
int16_t
__hsail_sat_add_s16 (int16_t a, int16_t b)
{
int32_t c = (int32_t) a + (int32_t) b;
if (c > INT16_MAX)
return INT16_MAX;
else if (c < INT16_MIN)
return INT16_MIN;
else
return c;
}
int32_t
__hsail_sat_add_s32 (int32_t a, int32_t b)
{
int32_t c;
if (__builtin_add_overflow (a, b, &c))
return b < 0 ? INT32_MIN : INT32_MAX;
return c;
}
int64_t
__hsail_sat_add_s64 (int64_t a, int64_t b)
{
int64_t c;
if (__builtin_add_overflow (a, b, &c))
return b < 0 ? INT64_MIN : INT64_MAX;
return c;
}
uint8_t
__hsail_sat_sub_u8 (uint8_t a, uint8_t b)
{
int16_t c = (uint16_t) a - (uint16_t) b;
if (c < 0)
return 0;
else
return c;
}
uint16_t
__hsail_sat_sub_u16 (uint16_t a, uint16_t b)
{
int32_t c = (uint32_t) a - (uint32_t) b;
if (c < 0)
return 0;
else
return c;
}
uint32_t
__hsail_sat_sub_u32 (uint32_t a, uint32_t b)
{
uint32_t c;
if (__builtin_sub_overflow (a, b, &c))
return 0;
return c;
}
uint64_t
__hsail_sat_sub_u64 (uint64_t a, uint64_t b)
{
uint64_t c;
if (__builtin_sub_overflow (a, b, &c))
return 0;
return c;
}
int8_t
__hsail_sat_sub_s8 (int8_t a, int8_t b)
{
int16_t c = (int16_t) a - (int16_t) b;
if (c > INT8_MAX)
return INT8_MAX;
else if (c < INT8_MIN)
return INT8_MIN;
else
return c;
}
int16_t
__hsail_sat_sub_s16 (int16_t a, int16_t b)
{
int32_t c = (int32_t) a - (int32_t) b;
if (c > INT16_MAX)
return INT16_MAX;
else if (c < INT16_MIN)
return INT16_MIN;
else
return c;
}
int32_t
__hsail_sat_sub_s32 (int32_t a, int32_t b)
{
int32_t c;
if (__builtin_sub_overflow (a, b, &c))
return b < 0 ? INT32_MAX : INT32_MIN;
return c;
}
int64_t
__hsail_sat_sub_s64 (int64_t a, int64_t b)
{
int64_t c;
if (__builtin_sub_overflow (a, b, &c))
return b < 0 ? INT64_MAX : INT64_MIN;
return c;
}
uint8_t
__hsail_sat_mul_u8 (uint8_t a, uint8_t b)
{
uint16_t c = (uint16_t) a * (uint16_t) b;
if (c > UINT8_MAX)
return UINT8_MAX;
else
return c;
}
uint16_t
__hsail_sat_mul_u16 (uint16_t a, uint16_t b)
{
uint32_t c = (uint32_t) a * (uint32_t) b;
if (c > UINT16_MAX)
return UINT16_MAX;
else
return c;
}
uint32_t
__hsail_sat_mul_u32 (uint32_t a, uint32_t b)
{
uint32_t c;
if (__builtin_mul_overflow (a, b, &c))
return UINT32_MAX;
return c;
}
uint64_t
__hsail_sat_mul_u64 (uint64_t a, uint64_t b)
{
uint64_t c;
if (__builtin_mul_overflow (a, b, &c))
return UINT64_MAX;
return c;
}
int8_t
__hsail_sat_mul_s8 (int8_t a, int8_t b)
{
int16_t c = (int16_t) a * (int16_t) b;
if (c > INT8_MAX)
return INT8_MAX;
else if (c < INT8_MIN)
return INT8_MIN;
else
return c;
}
int16_t
__hsail_sat_mul_s16 (int16_t a, int16_t b)
{
int32_t c = (int32_t) a * (int32_t) b;
if (c > INT16_MAX)
return INT16_MAX;
else if (c < INT16_MIN)
return INT16_MIN;
else
return c;
}
int32_t
__hsail_sat_mul_s32 (int32_t a, int32_t b)
{
int32_t c;
if (__builtin_mul_overflow (a, b, &c))
return ((a > 0) ^ (b > 0)) ? INT32_MIN : INT32_MAX;
return c;
}
int64_t
__hsail_sat_mul_s64 (int64_t a, int64_t b)
{
int64_t c;
if (__builtin_mul_overflow (a, b, &c))
return ((a > 0) ^ (b > 0)) ? INT64_MIN : INT64_MAX;
return c;
}

59
libhsail-rt/rt/segment.c
View File

@ -1,59 +0,0 @@
/* segment.c -- Builtins for HSAIL segment related instructions.
Copyright (C) 2015-2021 Free Software Foundation, Inc.
Contributed by Pekka Jaaskelainen <pekka.jaaskelainen@parmance.com>
for General Processor Tech.
Permission is hereby granted, free of charge, to any person obtaining a
copy of this software and associated documentation files
(the "Software"), to deal in the Software without restriction, including
without limitation the rights to use, copy, modify, merge, publish,
distribute, sublicense, and/or sell copies of the Software, and to
permit persons to whom the Software is furnished to do so, subject to
the following conditions:
The above copyright notice and this permission notice shall be included
in all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS
OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.
IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM,
DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR
OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE
USE OR OTHER DEALINGS IN THE SOFTWARE.
*/
#include "workitems.h"
uint32_t
__hsail_segmentp_private (uint64_t flat_addr, PHSAWorkItem *wi)
{
if (flat_addr == 0)
return 1;
else
return ((void *) (uintptr_t) flat_addr >= wi->wg->private_base_ptr
&& ((void *) (uintptr_t) flat_addr
< (wi->wg->private_base_ptr
+ wi->wg->private_segment_total_size)));
}
uint32_t
__hsail_segmentp_group (uint64_t flat_addr, PHSAWorkItem *wi)
{
if (flat_addr == 0)
return 1;
else
return ((void *) (uintptr_t) flat_addr >= wi->wg->group_base_ptr
&& ((void *) (uintptr_t) flat_addr
< (wi->wg->group_base_ptr
+ wi->launch_data->dp->group_segment_size)));
}
uint32_t
__hsail_segmentp_global (uint64_t flat_addr, PHSAWorkItem *wi)
{
return (flat_addr == 0
|| (!__hsail_segmentp_private (flat_addr, wi)
&& !__hsail_segmentp_group (flat_addr, wi)));
}

977
libhsail-rt/rt/workitems.c
View File

@ -1,977 +0,0 @@
/* workitems.c -- The main runtime entry that performs work-item execution in
various ways and the builtin functions closely related to the
implementation.
Copyright (C) 2015-2021 Free Software Foundation, Inc.
Contributed by Pekka Jaaskelainen <pekka.jaaskelainen@parmance.com>
for General Processor Tech.
Permission is hereby granted, free of charge, to any person obtaining a
copy of this software and associated documentation files
(the "Software"), to deal in the Software without restriction, including
without limitation the rights to use, copy, modify, merge, publish,
distribute, sublicense, and/or sell copies of the Software, and to
permit persons to whom the Software is furnished to do so, subject to
the following conditions:
The above copyright notice and this permission notice shall be included
in all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS
OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.
IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM,
DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR
OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE
USE OR OTHER DEALINGS IN THE SOFTWARE.
*/
/* The fiber based multiple work-item work-group execution uses ucontext
based user mode threading. However, if gccbrig is able to optimize the
kernel to a much faster work-group function that implements the multiple
WI execution using loops instead of fibers requiring slow context switches,
the fiber-based implementation won't be called.
*/
#include <stdlib.h>
#include <signal.h>
#include <string.h>
#include "workitems.h"
#include "phsa-rt.h"
#ifdef HAVE_FIBERS
#include "fibers.h"
#endif
#ifdef BENCHMARK_PHSA_RT
#include <stdio.h>
#include <time.h>
static uint64_t wi_count = 0;
static uint64_t wis_skipped = 0;
static uint64_t wi_total = 0;
static clock_t start_time;
#endif
#ifdef DEBUG_PHSA_RT
#include <stdio.h>
#endif
#define PRIVATE_SEGMENT_ALIGN 256
#define FIBER_STACK_SIZE (64*1024)
#define GROUP_SEGMENT_ALIGN 256
/* Preserve this amount of additional space in the alloca stack as we need to
store the alloca frame pointer to the alloca frame, thus must preserve
space for it. This thus supports at most 1024 functions with allocas in
a call chain. */
#define ALLOCA_OVERHEAD 1024*4
uint32_t __hsail_workitemabsid (uint32_t dim, PHSAWorkItem *context);
uint32_t __hsail_workitemid (uint32_t dim, PHSAWorkItem *context);
uint32_t __hsail_gridgroups (uint32_t dim, PHSAWorkItem *context);
uint32_t __hsail_currentworkgroupsize (uint32_t dim, PHSAWorkItem *wi);
uint32_t __hsail_workgroupsize (uint32_t dim, PHSAWorkItem *wi);
void
phsa_fatal_error (int code)
{
exit (code);
}
#ifdef HAVE_FIBERS
/* ucontext-based work-item thread implementation. Runs all work-items in
separate fibers. */
static void
phsa_work_item_thread (int arg0, int arg1)
{
void *arg = fiber_int_args_to_ptr (arg0, arg1);
PHSAWorkItem *wi = (PHSAWorkItem *) arg;
volatile PHSAWorkGroup *wg = wi->wg;
PHSAKernelLaunchData *l_data = wi->launch_data;
do
{
int retcode
= fiber_barrier_reach ((fiber_barrier_t *) l_data->wg_start_barrier);
/* At this point the threads can assume that either more_wgs is 0 or
the current_work_group_* is set to point to the WG executed next. */
if (!wi->wg->more_wgs)
break;
wi->group_x = wg->x;
wi->group_y = wg->y;
wi->group_z = wg->z;
wi->cur_wg_size_x = __hsail_currentworkgroupsize (0, wi);
wi->cur_wg_size_y = __hsail_currentworkgroupsize (1, wi);
wi->cur_wg_size_z = __hsail_currentworkgroupsize (2, wi);
#ifdef DEBUG_PHSA_RT
printf (
"Running work-item %lu/%lu/%lu for wg %lu/%lu/%lu / %lu/%lu/%lu...\n",
wi->x, wi->y, wi->z, wi->group_x, wi->group_y, wi->group_z,
l_data->wg_max_x, l_data->wg_max_y, l_data->wg_max_z);
#endif
if (wi->x < __hsail_currentworkgroupsize (0, wi)
&& wi->y < __hsail_currentworkgroupsize (1, wi)
&& wi->z < __hsail_currentworkgroupsize (2, wi))
{
l_data->kernel (l_data->kernarg_addr, wi, wg->group_base_ptr,
wg->initial_group_offset, wg->private_base_ptr);
#ifdef DEBUG_PHSA_RT
printf ("done.\n");
#endif
#ifdef BENCHMARK_PHSA_RT
wi_count++;
#endif
}
else
{
#ifdef DEBUG_PHSA_RT
printf ("skipped (partial WG).\n");
#endif
#ifdef BENCHMARK_PHSA_RT
wis_skipped++;
#endif
}
retcode
= fiber_barrier_reach ((fiber_barrier_t *)
l_data->wg_completion_barrier);
/* The first thread updates the WG to execute next etc. */
if (retcode == 0)
{
#ifdef EXECUTE_WGS_BACKWARDS
if (wg->x == l_data->wg_min_x)
{
wg->x = l_data->wg_max_x - 1;
if (wg->y == l_data->wg_min_y)
{
wg->y = l_data->wg_max_y - 1;
if (wg->z == l_data->wg_min_z)
wg->more_wgs = 0;
else
wg->z--;
}
else
wg->y--;
}
else
wg->x--;
#else
if (wg->x + 1 >= l_data->wg_max_x)
{
wg->x = l_data->wg_min_x;
if (wg->y + 1 >= l_data->wg_max_y)
{
wg->y = l_data->wg_min_y;
if (wg->z + 1 >= l_data->wg_max_z)
wg->more_wgs = 0;
else
wg->z++;
}
else
wg->y++;
}
else
wg->x++;
#endif
wi->group_x = wg->x;
wi->group_y = wg->y;
wi->group_z = wg->z;
wi->cur_wg_size_x = __hsail_currentworkgroupsize (0, wi);
wi->cur_wg_size_y = __hsail_currentworkgroupsize (1, wi);
wi->cur_wg_size_z = __hsail_currentworkgroupsize (2, wi);
/* Reinitialize the work-group barrier according to the new WG's
size, which might not be the same as the previous ones, due
to "partial WGs". */
size_t wg_size = __hsail_currentworkgroupsize (0, wi)
* __hsail_currentworkgroupsize (1, wi)
* __hsail_currentworkgroupsize (2, wi);
#ifdef DEBUG_PHSA_RT
printf ("Reinitializing the WG barrier to %lu.\n", wg_size);
#endif
fiber_barrier_init ((fiber_barrier_t *)
wi->launch_data->wg_sync_barrier,
wg_size);
#ifdef BENCHMARK_PHSA_RT
if (wi_count % 1000 == 0)
{
clock_t spent_time = clock () - start_time;
double spent_time_sec = (double) spent_time / CLOCKS_PER_SEC;
double wis_per_sec = wi_count / spent_time_sec;
uint64_t eta_sec
= (wi_total - wi_count - wis_skipped) / wis_per_sec;
printf ("%lu WIs executed %lu skipped in %lus (%lu WIs/s, ETA in "
"%lu s)\n",
wi_count, wis_skipped, (uint64_t) spent_time_sec,
(uint64_t) wis_per_sec, (uint64_t) eta_sec);
}
#endif
}
}
while (1);
fiber_exit ();
}
#endif
#define MIN(a, b) ((a < b) ? a : b)
#define MAX(a, b) ((a > b) ? a : b)
#ifdef HAVE_FIBERS
/* Spawns a given number of work-items to execute a set of work-groups,
blocks until their completion. */
static void
phsa_execute_wi_gang (PHSAKernelLaunchData *context, void *group_base_ptr,
uint32_t group_local_offset, size_t wg_size_x,
size_t wg_size_y, size_t wg_size_z)
{
PHSAWorkItem *wi_threads = NULL;
PHSAWorkGroup wg;
size_t flat_wi_id = 0, x, y, z, max_x, max_y, max_z;
uint32_t group_x, group_y, group_z;
fiber_barrier_t wg_start_barrier;
fiber_barrier_t wg_completion_barrier;
fiber_barrier_t wg_sync_barrier;
max_x = wg_size_x == 0 ? 1 : wg_size_x;
max_y = wg_size_y == 0 ? 1 : wg_size_y;
max_z = wg_size_z == 0 ? 1 : wg_size_z;
size_t wg_size = max_x * max_y * max_z;
if (wg_size > PHSA_MAX_WG_SIZE)
phsa_fatal_error (2);
wg.private_segment_total_size = context->dp->private_segment_size * wg_size;
if (wg.private_segment_total_size > 0
&& posix_memalign (&wg.private_base_ptr, PRIVATE_SEGMENT_ALIGN,
wg.private_segment_total_size)
!= 0)
phsa_fatal_error (3);
wg.alloca_stack_p = wg.private_segment_total_size + ALLOCA_OVERHEAD;
wg.alloca_frame_p = wg.alloca_stack_p;
wg.initial_group_offset = group_local_offset;
#ifdef EXECUTE_WGS_BACKWARDS
group_x = context->wg_max_x - 1;
group_y = context->wg_max_y - 1;
group_z = context->wg_max_z - 1;
#else
group_x = context->wg_min_x;
group_y = context->wg_min_y;
group_z = context->wg_min_z;
#endif
fiber_barrier_init (&wg_sync_barrier, wg_size);
fiber_barrier_init (&wg_start_barrier, wg_size);
fiber_barrier_init (&wg_completion_barrier, wg_size);
context->wg_start_barrier = &wg_start_barrier;
context->wg_sync_barrier = &wg_sync_barrier;
context->wg_completion_barrier = &wg_completion_barrier;
wg.more_wgs = 1;
wg.group_base_ptr = group_base_ptr;
#ifdef BENCHMARK_PHSA_RT
wi_count = 0;
wis_skipped = 0;
start_time = clock ();
#endif
wi_threads = malloc (sizeof (PHSAWorkItem) * max_x * max_y * max_z);
for (x = 0; x < max_x; ++x)
for (y = 0; y < max_y; ++y)
for (z = 0; z < max_z; ++z)
{
PHSAWorkItem *wi = &wi_threads[flat_wi_id];
wi->launch_data = context;
wi->wg = &wg;
wg.x = wi->group_x = group_x;
wg.y = wi->group_y = group_y;
wg.z = wi->group_z = group_z;
wi->wg_size_x = context->dp->workgroup_size_x;
wi->wg_size_y = context->dp->workgroup_size_y;
wi->wg_size_z = context->dp->workgroup_size_z;
wi->cur_wg_size_x = __hsail_currentworkgroupsize (0, wi);
wi->cur_wg_size_y = __hsail_currentworkgroupsize (1, wi);
wi->cur_wg_size_z = __hsail_currentworkgroupsize (2, wi);
wi->x = x;
wi->y = y;
wi->z = z;
/* TODO: set the stack size according to the private
segment size. Too big stack consumes huge amount of
memory in case of huge number of WIs and a too small stack
will fail in mysterious and potentially dangerous ways. */
fiber_init (&wi->fiber, phsa_work_item_thread, wi,
FIBER_STACK_SIZE, PRIVATE_SEGMENT_ALIGN);
++flat_wi_id;
}
do
{
--flat_wi_id;
fiber_join (&wi_threads[flat_wi_id].fiber);
}
while (flat_wi_id > 0);
if (wg.private_segment_total_size > 0)
free (wg.private_base_ptr);
free (wi_threads);
}
/* Spawn the work-item threads to execute work-groups and let
them execute all the WGs, including a potential partial WG. */
static void
phsa_spawn_work_items (PHSAKernelLaunchData *context, void *group_base_ptr,
uint32_t group_local_offset)
{
hsa_kernel_dispatch_packet_t *dp = context->dp;
size_t x, y, z;
context->group_segment_start_addr = (size_t) group_base_ptr;
/* HSA seems to allow the WG size to be larger than the grid size. We need to
saturate the effective WG size to the grid size to prevent the extra WIs
from executing. */
size_t sat_wg_size_x, sat_wg_size_y, sat_wg_size_z, sat_wg_size;
sat_wg_size_x = MIN (dp->workgroup_size_x, dp->grid_size_x);
sat_wg_size_y = MIN (dp->workgroup_size_y, dp->grid_size_y);
sat_wg_size_z = MIN (dp->workgroup_size_z, dp->grid_size_z);
sat_wg_size = sat_wg_size_x * sat_wg_size_y * sat_wg_size_z;
#ifdef BENCHMARK_PHSA_RT
wi_total = (uint64_t) dp->grid_size_x
* (dp->grid_size_y > 0 ? dp->grid_size_y : 1)
* (dp->grid_size_z > 0 ? dp->grid_size_z : 1);
#endif
/* For now execute all work groups in a single coarse thread (does not utilize
multicore/multithread). */
context->wg_min_x = context->wg_min_y = context->wg_min_z = 0;
int dims = dp->setup & 0x3;
context->wg_max_x = ((uint64_t) dp->grid_size_x + dp->workgroup_size_x - 1)
/ dp->workgroup_size_x;
context->wg_max_y
= dims < 2 ? 1 : ((uint64_t) dp->grid_size_y + dp->workgroup_size_y - 1)
/ dp->workgroup_size_y;
context->wg_max_z
= dims < 3 ? 1 : ((uint64_t) dp->grid_size_z + dp->workgroup_size_z - 1)
/ dp->workgroup_size_z;
#ifdef DEBUG_PHSA_RT
printf ("### launching work-groups %lu/%lu/%lu to %lu/%lu/%lu with "
"wg size %lu/%lu/%lu grid size %u/%u/%u\n",
context->wg_min_x, context->wg_min_y, context->wg_min_z,
context->wg_max_x, context->wg_max_y, context->wg_max_z,
sat_wg_size_x, sat_wg_size_y, sat_wg_size_z, dp->grid_size_x,
dp->grid_size_y, dp->grid_size_z);
#endif
phsa_execute_wi_gang (context, group_base_ptr, group_local_offset,
sat_wg_size_x, sat_wg_size_y, sat_wg_size_z);
}
#endif
/* Executes the given work-group function for all work groups in the grid.
A work-group function is a version of the original kernel which executes
the kernel for all work-items in a work-group. It is produced by gccbrig
if it can handle the kernel's barrier usage and is much faster way to
execute massive numbers of work-items in a non-SPMD machine than fibers
(easily 100x faster). */
static void
phsa_execute_work_groups (PHSAKernelLaunchData *context, void *group_base_ptr,
uint32_t group_local_offset)
{
hsa_kernel_dispatch_packet_t *dp = context->dp;
size_t x, y, z, wg_x, wg_y, wg_z;
context->group_segment_start_addr = (size_t) group_base_ptr;
/* HSA seems to allow the WG size to be larger than the grid size. We need
to saturate the effective WG size to the grid size to prevent the extra WIs
from executing. */
size_t sat_wg_size_x, sat_wg_size_y, sat_wg_size_z, sat_wg_size;
sat_wg_size_x = MIN (dp->workgroup_size_x, dp->grid_size_x);
sat_wg_size_y = MIN (dp->workgroup_size_y, dp->grid_size_y);
sat_wg_size_z = MIN (dp->workgroup_size_z, dp->grid_size_z);
sat_wg_size = sat_wg_size_x * sat_wg_size_y * sat_wg_size_z;
#ifdef BENCHMARK_PHSA_RT
wi_total = (uint64_t) dp->grid_size_x
* (dp->grid_size_y > 0 ? dp->grid_size_y : 1)
* (dp->grid_size_z > 0 ? dp->grid_size_z : 1);
#endif
context->wg_min_x = context->wg_min_y = context->wg_min_z = 0;
int dims = dp->setup & 0x3;
context->wg_max_x = ((uint64_t) dp->grid_size_x + dp->workgroup_size_x - 1)
/ dp->workgroup_size_x;
context->wg_max_y
= dims < 2 ? 1 : ((uint64_t) dp->grid_size_y + dp->workgroup_size_y - 1)
/ dp->workgroup_size_y;
context->wg_max_z
= dims < 3 ? 1 : ((uint64_t) dp->grid_size_z + dp->workgroup_size_z - 1)
/ dp->workgroup_size_z;
#ifdef DEBUG_PHSA_RT
printf ("### launching work-groups %lu/%lu/%lu to %lu/%lu/%lu with "
"wg size %lu/%lu/%lu grid size %u/%u/%u\n",
context->wg_min_x, context->wg_min_y, context->wg_min_z,
context->wg_max_x, context->wg_max_y, context->wg_max_z,
sat_wg_size_x, sat_wg_size_y, sat_wg_size_z, dp->grid_size_x,
dp->grid_size_y, dp->grid_size_z);
#endif
PHSAWorkItem wi;
PHSAWorkGroup wg;
wi.wg = &wg;
wi.x = wi.y = wi.z = 0;
wi.launch_data = context;
#ifdef BENCHMARK_PHSA_RT
start_time = clock ();
uint64_t wg_count = 0;
#endif
size_t wg_size = __hsail_workgroupsize (0, &wi)
* __hsail_workgroupsize (1, &wi)
* __hsail_workgroupsize (2, &wi);
void *private_base_ptr = NULL;
if (dp->private_segment_size > 0
&& posix_memalign (&private_base_ptr, PRIVATE_SEGMENT_ALIGN,
dp->private_segment_size * wg_size)
!= 0)
phsa_fatal_error (3);
wg.alloca_stack_p = dp->private_segment_size * wg_size + ALLOCA_OVERHEAD;
wg.alloca_frame_p = wg.alloca_stack_p;
wg.private_base_ptr = private_base_ptr;
wg.group_base_ptr = group_base_ptr;
#ifdef DEBUG_PHSA_RT
printf ("priv seg size %u wg_size %lu @ %p\n", dp->private_segment_size,
wg_size, private_base_ptr);
#endif
for (wg_z = context->wg_min_z; wg_z < context->wg_max_z; ++wg_z)
for (wg_y = context->wg_min_y; wg_y < context->wg_max_y; ++wg_y)
for (wg_x = context->wg_min_x; wg_x < context->wg_max_x; ++wg_x)
{
wi.group_x = wg_x;
wi.group_y = wg_y;
wi.group_z = wg_z;
wi.wg_size_x = context->dp->workgroup_size_x;
wi.wg_size_y = context->dp->workgroup_size_y;
wi.wg_size_z = context->dp->workgroup_size_z;
wi.cur_wg_size_x = __hsail_currentworkgroupsize (0, &wi);
wi.cur_wg_size_y = __hsail_currentworkgroupsize (1, &wi);
wi.cur_wg_size_z = __hsail_currentworkgroupsize (2, &wi);
context->kernel (context->kernarg_addr, &wi, group_base_ptr,
group_local_offset, private_base_ptr);
#if defined (BENCHMARK_PHSA_RT)
wg_count++;
if (wg_count % 1000000 == 0)
{
clock_t spent_time = clock () - start_time;
uint64_t wi_count = wg_x * sat_wg_size_x + wg_y * sat_wg_size_y
+ wg_z * sat_wg_size_z;
double spent_time_sec = (double) spent_time / CLOCKS_PER_SEC;
double wis_per_sec = wi_count / spent_time_sec;
uint64_t eta_sec = (wi_total - wi_count) / wis_per_sec;
printf ("%lu WIs executed in %lus (%lu WIs/s, ETA in %lu s)\n",
wi_count, (uint64_t) spent_time_sec,
(uint64_t) wis_per_sec, (uint64_t) eta_sec);
}
#endif
}
#ifdef BENCHMARK_PHSA_RT
clock_t spent_time = clock () - start_time;
double spent_time_sec = (double) spent_time / CLOCKS_PER_SEC;
double wis_per_sec = wi_total / spent_time_sec;
printf ("### %lu WIs executed in %lu s (%lu WIs / s)\n", wi_total,
(uint64_t) spent_time_sec, (uint64_t) wis_per_sec);
#endif
free (private_base_ptr);
private_base_ptr = NULL;
}
/* gccbrig generates the following from each HSAIL kernel:
1) The actual kernel function (a single work-item kernel or a work-group
function) generated from HSAIL (BRIG).
static void _Kernel (void* args, void* context, void* group_base_ptr)
{
...
}
2) A public facing kernel function that is called from the PHSA runtime:
a) A single work-item function (that requires fibers for multi-WI):
void Kernel (void* context)
{
__launch_launch_kernel (_Kernel, context);
}
or
b) a when gccbrig could generate a work-group function:
void Kernel (void* context)
{
__hsail_launch_wg_function (_Kernel, context);
}
*/
#ifdef HAVE_FIBERS
void
__hsail_launch_kernel (gccbrigKernelFunc kernel, PHSAKernelLaunchData *context,
void *group_base_ptr, uint32_t group_local_offset)
{
context->kernel = kernel;
phsa_spawn_work_items (context, group_base_ptr, group_local_offset);
}
#endif
void
__hsail_launch_wg_function (gccbrigKernelFunc kernel,
PHSAKernelLaunchData *context, void *group_base_ptr,
uint32_t group_local_offset)
{
context->kernel = kernel;
phsa_execute_work_groups (context, group_base_ptr, group_local_offset);
}
uint32_t
__hsail_workitemabsid (uint32_t dim, PHSAWorkItem *context)
{
hsa_kernel_dispatch_packet_t *dp = context->launch_data->dp;
uint32_t id;
switch (dim)
{
default:
case 0:
/* Overflow semantics in the case of WG dim > grid dim. */
id = ((uint64_t) context->group_x * dp->workgroup_size_x + context->x)
% dp->grid_size_x;
break;
case 1:
id = ((uint64_t) context->group_y * dp->workgroup_size_y + context->y)
% dp->grid_size_y;
break;
case 2:
id = ((uint64_t) context->group_z * dp->workgroup_size_z + context->z)
% dp->grid_size_z;
break;
}
return id;
}
uint64_t
__hsail_workitemabsid_u64 (uint32_t dim, PHSAWorkItem *context)
{
hsa_kernel_dispatch_packet_t *dp = context->launch_data->dp;
uint64_t id;
switch (dim)
{
default:
case 0:
/* Overflow semantics in the case of WG dim > grid dim. */
id = ((uint64_t) context->group_x * dp->workgroup_size_x + context->x)
% dp->grid_size_x;
break;
case 1:
id = ((uint64_t) context->group_y * dp->workgroup_size_y + context->y)
% dp->grid_size_y;
break;
case 2:
id = ((uint64_t) context->group_z * dp->workgroup_size_z + context->z)
% dp->grid_size_z;
break;
}
return id;
}
uint32_t
__hsail_workitemid (uint32_t dim, PHSAWorkItem *context)
{
PHSAWorkItem *c = (PHSAWorkItem *) context;
hsa_kernel_dispatch_packet_t *dp = context->launch_data->dp;
/* The number of dimensions is in the two least significant bits. */
int dims = dp->setup & 0x3;
uint32_t id;
switch (dim)
{
default:
case 0:
id = c->x;
break;
case 1:
id = dims < 2 ? 0 : c->y;
break;
case 2:
id = dims < 3 ? 0 : c->z;
break;
}
return id;
}
uint32_t
__hsail_gridgroups (uint32_t dim, PHSAWorkItem *context)
{
hsa_kernel_dispatch_packet_t *dp = context->launch_data->dp;
int dims = dp->setup & 0x3;
uint32_t id;
switch (dim)
{
default:
case 0:
id = (dp->grid_size_x + dp->workgroup_size_x - 1) / dp->workgroup_size_x;
break;
case 1:
id = dims < 2 ? 1 : (dp->grid_size_y + dp->workgroup_size_y - 1)
/ dp->workgroup_size_y;
break;
case 2:
id = dims < 3 ? 1 : (dp->grid_size_z + dp->workgroup_size_z - 1)
/ dp->workgroup_size_z;
break;
}
return id;
}
uint32_t
__hsail_workitemflatid (PHSAWorkItem *c)
{
hsa_kernel_dispatch_packet_t *dp = c->launch_data->dp;
return c->x + c->y * dp->workgroup_size_x
+ c->z * dp->workgroup_size_x * dp->workgroup_size_y;
}
uint32_t
__hsail_currentworkitemflatid (PHSAWorkItem *c)
{
hsa_kernel_dispatch_packet_t *dp = c->launch_data->dp;
return c->x + c->y * __hsail_currentworkgroupsize (0, c)
+ c->z * __hsail_currentworkgroupsize (0, c)
* __hsail_currentworkgroupsize (1, c);
}
void
__hsail_setworkitemid (uint32_t dim, uint32_t id, PHSAWorkItem *context)
{
switch (dim)
{
default:
case 0:
context->x = id;
break;
case 1:
context->y = id;
break;
case 2:
context->z = id;
break;
}
}
uint64_t
__hsail_workitemflatabsid_u64 (PHSAWorkItem *context)
{
PHSAWorkItem *c = (PHSAWorkItem *) context;
hsa_kernel_dispatch_packet_t *dp = context->launch_data->dp;
/* Work-item flattened absolute ID = ID0 + ID1 * max0 + ID2 * max0 * max1. */
uint64_t id0 = __hsail_workitemabsid (0, context);
uint64_t id1 = __hsail_workitemabsid (1, context);
uint64_t id2 = __hsail_workitemabsid (2, context);
uint64_t max0 = dp->grid_size_x;
uint64_t max1 = dp->grid_size_y;
uint64_t id = id0 + id1 * max0 + id2 * max0 * max1;
return id;
}
uint32_t
__hsail_workitemflatabsid_u32 (PHSAWorkItem *context)
{
PHSAWorkItem *c = (PHSAWorkItem *) context;
hsa_kernel_dispatch_packet_t *dp = context->launch_data->dp;
/* work-item flattened absolute ID = ID0 + ID1 * max0 + ID2 * max0 * max1. */
uint64_t id0 = __hsail_workitemabsid (0, context);
uint64_t id1 = __hsail_workitemabsid (1, context);
uint64_t id2 = __hsail_workitemabsid (2, context);
uint64_t max0 = dp->grid_size_x;
uint64_t max1 = dp->grid_size_y;
uint64_t id = id0 + id1 * max0 + id2 * max0 * max1;
return (uint32_t) id;
}
uint32_t
__hsail_currentworkgroupsize (uint32_t dim, PHSAWorkItem *wi)
{
hsa_kernel_dispatch_packet_t *dp = wi->launch_data->dp;
uint32_t wg_size = 0;
switch (dim)
{
default:
case 0:
if ((uint64_t) wi->group_x < dp->grid_size_x / dp->workgroup_size_x)
wg_size = dp->workgroup_size_x; /* Full WG. */
else
wg_size = dp->grid_size_x % dp->workgroup_size_x; /* Partial WG. */
break;
case 1:
if ((uint64_t) wi->group_y < dp->grid_size_y / dp->workgroup_size_y)
wg_size = dp->workgroup_size_y; /* Full WG. */
else
wg_size = dp->grid_size_y % dp->workgroup_size_y; /* Partial WG. */
break;
case 2:
if ((uint64_t) wi->group_z < dp->grid_size_z / dp->workgroup_size_z)
wg_size = dp->workgroup_size_z; /* Full WG. */
else
wg_size = dp->grid_size_z % dp->workgroup_size_z; /* Partial WG. */
break;
}
return wg_size;
}
uint32_t
__hsail_workgroupsize (uint32_t dim, PHSAWorkItem *wi)
{
hsa_kernel_dispatch_packet_t *dp = wi->launch_data->dp;
switch (dim)
{
default:
case 0:
return dp->workgroup_size_x;
case 1:
return dp->workgroup_size_y;
case 2:
return dp->workgroup_size_z;
}
}
uint32_t
__hsail_gridsize (uint32_t dim, PHSAWorkItem *wi)
{
hsa_kernel_dispatch_packet_t *dp = wi->launch_data->dp;
switch (dim)
{
default:
case 0:
return dp->grid_size_x;
case 1:
return dp->grid_size_y;
case 2:
return dp->grid_size_z;
}
}
uint32_t
__hsail_workgroupid (uint32_t dim, PHSAWorkItem *wi)
{
switch (dim)
{
default:
case 0:
return wi->group_x;
case 1:
return wi->group_y;
case 2:
return wi->group_z;
}
}
uint32_t
__hsail_dim (PHSAWorkItem *wi)
{
hsa_kernel_dispatch_packet_t *dp = wi->launch_data->dp;
return dp->setup & 0x3;
}
uint64_t
__hsail_packetid (PHSAWorkItem *wi)
{
return wi->launch_data->packet_id;
}
uint32_t
__hsail_packetcompletionsig_sig32 (PHSAWorkItem *wi)
{
return (uint32_t) wi->launch_data->dp->completion_signal.handle;
}
uint64_t
__hsail_packetcompletionsig_sig64 (PHSAWorkItem *wi)
{
return (uint64_t) (wi->launch_data->dp->completion_signal.handle);
}
#ifdef HAVE_FIBERS
void
__hsail_barrier (PHSAWorkItem *wi)
{
fiber_barrier_reach ((fiber_barrier_t *) wi->launch_data->wg_sync_barrier);
}
#endif
/* Return a 32b private segment address that points to a dynamically
allocated chunk of 'size' with 'align'.
Allocates the space from the end of the private segment allocated
for the whole work group. In implementations with separate private
memories per WI, we will need to have a stack pointer per WI. But in
the current implementation, the segment is shared, so we possibly
save some space in case all WIs do not call the alloca.
The "alloca frames" are organized as follows:
wg->alloca_stack_p points to the last allocated data (initially
outside the private segment)
wg->alloca_frame_p points to the first address _outside_ the current
function's allocations (initially to the same as alloca_stack_p)
The data is allocated downwards from the end of the private segment.
In the beginning of a new function which has allocas, a new alloca
frame is pushed which adds the current alloca_frame_p (the current
function's frame starting point) to the top of the alloca stack and
alloca_frame_p is set to the current stack position.
At the exit points of a function with allocas, the alloca frame
is popped before returning. This involves popping the alloca_frame_p
to the one of the previous function in the call stack, and alloca_stack_p
similarly, to the position of the last word alloca'd by the previous
function.
*/
uint32_t
__hsail_alloca (uint32_t size, uint32_t align, PHSAWorkItem *wi)
{
volatile PHSAWorkGroup *wg = wi->wg;
int64_t new_pos = wg->alloca_stack_p - size;
while (new_pos % align != 0)
new_pos--;
if (new_pos < 0)
phsa_fatal_error (2);
wg->alloca_stack_p = new_pos;
#ifdef DEBUG_ALLOCA
printf ("--- alloca (%u, %u) sp @%u fp @%u\n", size, align,
wg->alloca_stack_p, wg->alloca_frame_p);
#endif
return new_pos;
}
/* Initializes a new "alloca frame" in the private segment.
This should be called at all the function entry points in case
the function contains at least one call to alloca. */
void
__hsail_alloca_push_frame (PHSAWorkItem *wi)
{
volatile PHSAWorkGroup *wg = wi->wg;
/* Store the alloca_frame_p without any alignment padding so
we know exactly where the previous frame ended after popping
it. */
#ifdef DEBUG_ALLOCA
printf ("--- push frame ");
#endif
uint32_t last_word_offs = __hsail_alloca (4, 1, wi);
memcpy (wg->private_base_ptr + last_word_offs,
(const void *) &wg->alloca_frame_p, 4);
wg->alloca_frame_p = last_word_offs;
#ifdef DEBUG_ALLOCA
printf ("--- sp @%u fp @%u\n", wg->alloca_stack_p, wg->alloca_frame_p);
#endif
}
/* Frees the current "alloca frame" and restores the frame
pointer.
This should be called at all the function return points in case
the function contains at least one call to alloca. Restores the
alloca stack to the condition it was before pushing the frame
the last time. */
void
__hsail_alloca_pop_frame (PHSAWorkItem *wi)
{
volatile PHSAWorkGroup *wg = wi->wg;
wg->alloca_stack_p = wg->alloca_frame_p;
memcpy ((void *) &wg->alloca_frame_p,
(const void *) (wg->private_base_ptr + wg->alloca_frame_p), 4);
/* Now frame_p points to the beginning of the previous function's
frame and stack_p to its end. */
wg->alloca_stack_p += 4;
#ifdef DEBUG_ALLOCA
printf ("--- pop frame sp @%u fp @%u\n", wg->alloca_stack_p,
wg->alloca_frame_p);
#endif
}

68
libhsail-rt/target-config.h.in
View File

@ -1,68 +0,0 @@
/* target-config.h.in. Generated from configure.ac by autoheader. */
/* Define to 1 if you have the <dlfcn.h> header file. */
#undef HAVE_DLFCN_H
/* Define to 1 if you have the <inttypes.h> header file. */
#undef HAVE_INTTYPES_H
/* Define to 1 if you have the <memory.h> header file. */
#undef HAVE_MEMORY_H
/* Define to 1 if you have the <stdint.h> header file. */
#undef HAVE_STDINT_H
/* Define to 1 if you have the <stdlib.h> header file. */
#undef HAVE_STDLIB_H
/* Define to 1 if you have the <strings.h> header file. */
#undef HAVE_STRINGS_H
/* Define to 1 if you have the <string.h> header file. */
#undef HAVE_STRING_H
/* Define to 1 if you have the <sys/stat.h> header file. */
#undef HAVE_SYS_STAT_H
/* Define to 1 if you have the <sys/types.h> header file. */
#undef HAVE_SYS_TYPES_H
/* Define to 1 if you have the <unistd.h> header file. */
#undef HAVE_UNISTD_H
/* Define to the sub-directory in which libtool stores uninstalled libraries.
*/
#undef LT_OBJDIR
/* Name of package */
#undef PACKAGE
/* Define to the address where bug reports for this package should be sent. */
#undef PACKAGE_BUGREPORT
/* Define to the full name of this package. */
#undef PACKAGE_NAME
/* Define to the full name and version of this package. */
#undef PACKAGE_STRING
/* Define to the one symbol short name of this package. */
#undef PACKAGE_TARNAME
/* Define to the home page for this package. */
#undef PACKAGE_URL
/* Define to the version of this package. */
#undef PACKAGE_VERSION
/* The size of `int', as computed by sizeof. */
#undef SIZEOF_INT
/* The size of `void*', as computed by sizeof. */
#undef SIZEOF_VOIDP
/* Define to 1 if you have the ANSI C header files. */
#undef STDC_HEADERS
/* Version number of package */
#undef VERSION