Brig front-end

2017-01-24  Pekka Jääskeläinen <pekka@parmance.com>
	    Martin Jambor  <mjambor@suse.cz>

	* Makefile.def (target_modules): Added libhsail-rt.
	(languages): Added language brig.
	* Makefile.in: Regenerated.
	* configure.ac (TOPLEVEL_CONFIGURE_ARGUMENTS): Added
	tgarget-libhsail-rt.  Make brig unsupported on untested architectures.
	* configure: Regenerated.

gcc/
	* brig-builtins.def: New file.
	* builtins.def (DEF_HSAIL_BUILTIN): New macro.
	(DEF_HSAIL_ATOMIC_BUILTIN): Likewise.
	(DEF_HSAIL_SAT_BUILTIN): Likewise.
	(DEF_HSAIL_INTR_BUILTIN): Likewise.
	(DEF_HSAIL_CVT_ZEROI_SAT_BUILTIN): Likewise.
	* builtin-types.def (BT_INT8): New.
	(BT_INT16): Likewise.
	(BT_UINT8): Likewise.
	(BT_UINT16): Likewise.
	(BT_FN_ULONG): Likewise.
	(BT_FN_UINT_INT): Likewise.
	(BT_FN_UINT_ULONG): Likewise.
	(BT_FN_UINT_LONG): Likewise.
	(BT_FN_UINT_PTR): Likewise.
	(BT_FN_ULONG_PTR): Likewise.
	(BT_FN_INT8_FLOAT): Likewise.
	(BT_FN_INT16_FLOAT): Likewise.
	(BT_FN_UINT32_FLOAT): Likewise.
	(BT_FN_UINT16_FLOAT): Likewise.
	(BT_FN_UINT8_FLOAT): Likewise.
	(BT_FN_UINT64_FLOAT): Likewise.
	(BT_FN_UINT16_UINT32): Likewise.
	(BT_FN_UINT32_UINT16): Likewise.
	(BT_FN_UINT16_UINT16_UINT16): Likewise.
	(BT_FN_INT_PTR_INT): Likewise.
	(BT_FN_UINT_PTR_UINT): Likewise.
	(BT_FN_LONG_PTR_LONG): Likewise.
	(BT_FN_ULONG_PTR_ULONG): Likewise.
	(BT_FN_VOID_UINT64_UINT64): Likewise.
	(BT_FN_UINT8_UINT8_UINT8): Likewise.
	(BT_FN_INT8_INT8_INT8): Likewise.
	(BT_FN_INT16_INT16_INT16): Likewise.
	(BT_FN_INT_INT_INT): Likewise.
	(BT_FN_UINT_FLOAT_UINT): Likewise.
	(BT_FN_FLOAT_UINT_UINT): Likewise.
	(BT_FN_ULONG_UINT_UINT): Likewise.
	(BT_FN_ULONG_UINT_PTR): Likewise.
	(BT_FN_ULONG_ULONG_ULONG): Likewise.
	(BT_FN_UINT_UINT_UINT): Likewise.
	(BT_FN_VOID_UINT_PTR): Likewise.
	(BT_FN_UINT_UINT_PTR: Likewise.
	(BT_FN_UINT32_UINT64_PTR): Likewise.
	(BT_FN_INT_INT_UINT_UINT): Likewise.
	(BT_FN_UINT_UINT_UINT_UINT): Likewise.
	(BT_FN_UINT_UINT_UINT_PTR): Likewise.
	(BT_FN_UINT_ULONG_ULONG_UINT): Likewise.
	(BT_FN_ULONG_ULONG_ULONG_ULONG): Likewise.
	(BT_FN_LONG_LONG_UINT_UINT): Likewise.
	(BT_FN_ULONG_ULONG_UINT_UINT): Likewise.
	(BT_FN_VOID_UINT32_UINT64_PTR): Likewise.
	(BT_FN_VOID_UINT32_UINT32_PTR): Likewise.
	(BT_FN_UINT_UINT_UINT_UINT_UINT): Likewise.
	(BT_FN_UINT_FLOAT_FLOAT_FLOAT_FLOAT): Likewise.
	(BT_FN_ULONG_ULONG_ULONG_UINT_UINT): Likewise.
	* doc/frontends.texi: List BRIG FE.
	* doc/install.texi (Testing): Add BRIG tesring requirements.
	* doc/invoke.texi (Overall Options): Mention BRIG.
	* doc/standards.texi (Standards): Doucment BRIG HSA version.

gcc/brig/

	* Make-lang.in: New file.
	* brig-builtins.h: Likewise.
	* brig-c.h: Likewise.
	* brig-lang.c: Likewise.
	* brigspec.c: Likewise.
	* config-lang.in: Likewise.
	* lang-specs.h: Likewise.
	* lang.opt: Likewise.
	* brigfrontend/brig-arg-block-handler.cc: Likewise.
	* brigfrontend/brig-atomic-inst-handler.cc: Likewise.
	* brigfrontend/brig-basic-inst-handler.cc: Likewise.
	* brigfrontend/brig-branch-inst-handler.cc: Likewise.
	* brigfrontend/brig-cmp-inst-handler.cc: Likewise.
	* brigfrontend/brig-code-entry-handler.cc: Likewise.
	* brigfrontend/brig-code-entry-handler.h: Likewise.
	* brigfrontend/brig-comment-handler.cc: Likewise.
	* brigfrontend/brig-control-handler.cc: Likewise.
	* brigfrontend/brig-copy-move-inst-handler.cc: Likewise.
	* brigfrontend/brig-cvt-inst-handler.cc: Likewise.
	* brigfrontend/brig-fbarrier-handler.cc: Likewise.
	* brigfrontend/brig-function-handler.cc: Likewise.
	* brigfrontend/brig-function.cc: Likewise.
	* brigfrontend/brig-function.h: Likewise.
	* brigfrontend/brig-inst-mod-handler.cc: Likewise.
	* brigfrontend/brig-label-handler.cc: Likewise.
	* brigfrontend/brig-lane-inst-handler.cc: Likewise.
	* brigfrontend/brig-machine.c: Likewise.
	* brigfrontend/brig-machine.h: Likewise.
	* brigfrontend/brig-mem-inst-handler.cc: Likewise.
	* brigfrontend/brig-module-handler.cc: Likewise.
	* brigfrontend/brig-queue-inst-handler.cc: Likewise.
	* brigfrontend/brig-seg-inst-handler.cc: Likewise.
	* brigfrontend/brig-signal-inst-handler.cc: Likewise.
	* brigfrontend/brig-to-generic.cc: Likewise.
	* brigfrontend/brig-to-generic.h: Likewise.
	* brigfrontend/brig-util.cc: Likewise.
	* brigfrontend/brig-util.h: Likewise.
	* brigfrontend/brig-variable-handler.cc: Likewise.
	* brigfrontend/phsa.h: Likewise.


gcc/testsuite/

	* lib/brig-dg.exp: New file.
	* lib/brig.exp: Likewise.
	* brig.dg/README: Likewise.
	* brig.dg/dg.exp: Likewise.
	* brig.dg/test/gimple/alloca.hsail: Likewise.
	* brig.dg/test/gimple/atomics.hsail: Likewise.
	* brig.dg/test/gimple/branches.hsail: Likewise.
	* brig.dg/test/gimple/fbarrier.hsail: Likewise.
	* brig.dg/test/gimple/function_calls.hsail: Likewise.
	* brig.dg/test/gimple/kernarg.hsail: Likewise.
	* brig.dg/test/gimple/mem.hsail: Likewise.
	* brig.dg/test/gimple/mulhi.hsail: Likewise.
	* brig.dg/test/gimple/packed.hsail: Likewise.
	* brig.dg/test/gimple/smoke_test.hsail: Likewise.
	* brig.dg/test/gimple/variables.hsail: Likewise.
	* brig.dg/test/gimple/vector.hsail: Likewise.

include/

	* hsa.h: Moved here from libgomp/plugin/hsa.h.

libgomp/

	* plugin/hsa.h: Moved to top level include.
	* plugin/plugin-hsa.c: Chanfgd include of hsa.h accordingly.

libhsail-rt/

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


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

From-SVN: r244867

ChangeLog

+2017-01-24  Pekka Jääskeläinen <pekka@parmance.com>
+	    Martin Jambor  <mjambor@suse.cz>
+
+
+	* Makefile.def (target_modules): Added libhsail-rt.
+	(languages): Added language brig.
+	* Makefile.in: Regenerated.
+	* configure.ac (TOPLEVEL_CONFIGURE_ARGUMENTS): Added
+	tgarget-libhsail-rt.  Make brig unsupported on untested architectures.
+	* configure: Regenerated.
+
 2017-01-19  Uros Bizjak  <ubizjak@gmail.com>
 
 	PR target/78478

Makefile.def

@@ -157,6 +157,7 @@ target_modules = { module= libquadmath; };
 target_modules = { module= libgfortran; };
 target_modules = { module= libobjc; };
 target_modules = { module= libgo; };
+target_modules = { module= libhsail-rt; };
 target_modules = { module= libtermcap; no_check=true;
                    missing=mostlyclean;
                    missing=clean;
@@ -601,6 +602,8 @@ languages = { language=objc;	gcc-check-target=check-objc;
 languages = { language=obj-c++;	gcc-check-target=check-obj-c++; };
 languages = { language=go;	gcc-check-target=check-go;
 				lib-check-target=check-target-libgo; };
+languages = { language=brig;	gcc-check-target=check-brig;
+				lib-check-target=check-target-libhsail-rt; };
 
 // Toplevel bootstrap
 bootstrap_stage = { id=1 ; };

configure

@@ -2754,6 +2754,7 @@ target_libraries="target-libgcc \
 		target-libgomp \
 		target-libcilkrts \
 		target-liboffloadmic \
+		target-libhsail-rt \
 		target-libatomic \
 		target-libitm \
 		target-libstdc++-v3 \
@@ -3482,6 +3483,19 @@ if test x$enable_libgo = x; then
     esac
 fi
 
+# 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*)
+    ;;
+    *)
+    unsupported_languages="$unsupported_languages brig"
+    # This implicitly disables also target-libhsail-rt as it won't
+    # get added to the build without BRIG FE.
+    ;;
+esac
+
 # Default libgloss CPU subdirectory.
 libgloss_dir="$target_cpu"
 

configure.ac

@@ -152,6 +152,7 @@ target_libraries="target-libgcc \
 		target-libgomp \
 		target-libcilkrts \
 		target-liboffloadmic \
+		target-libhsail-rt \
 		target-libatomic \
 		target-libitm \
 		target-libstdc++-v3 \
@@ -812,6 +813,19 @@ if test x$enable_libgo = x; then
     esac
 fi
 
+# 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*)
+    ;;
+    *)
+    unsupported_languages="$unsupported_languages brig"
+    # This implicitly disables also target-libhsail-rt as it won't
+    # get added to the build without BRIG FE.
+    ;;
+esac
+
 # Default libgloss CPU subdirectory.
 libgloss_dir="$target_cpu"
 

gcc/ChangeLog

+2017-01-24  Pekka Jääskeläinen <pekka@parmance.com>
+	    Martin Jambor  <mjambor@suse.cz>
+
+	* brig-builtins.def: New file.
+	* builtins.def (DEF_HSAIL_BUILTIN): New macro.
+	(DEF_HSAIL_ATOMIC_BUILTIN): Likewise.
+	(DEF_HSAIL_SAT_BUILTIN): Likewise.
+	(DEF_HSAIL_INTR_BUILTIN): Likewise.
+	(DEF_HSAIL_CVT_ZEROI_SAT_BUILTIN): Likewise.
+	* builtin-types.def (BT_INT8): New.
+	(BT_INT16): Likewise.
+	(BT_UINT8): Likewise.
+	(BT_UINT16): Likewise.
+	(BT_FN_ULONG): Likewise.
+	(BT_FN_UINT_INT): Likewise.
+	(BT_FN_UINT_ULONG): Likewise.
+	(BT_FN_UINT_LONG): Likewise.
+	(BT_FN_UINT_PTR): Likewise.
+	(BT_FN_ULONG_PTR): Likewise.
+	(BT_FN_INT8_FLOAT): Likewise.
+	(BT_FN_INT16_FLOAT): Likewise.
+	(BT_FN_UINT32_FLOAT): Likewise.
+	(BT_FN_UINT16_FLOAT): Likewise.
+	(BT_FN_UINT8_FLOAT): Likewise.
+	(BT_FN_UINT64_FLOAT): Likewise.
+	(BT_FN_UINT16_UINT32): Likewise.
+	(BT_FN_UINT32_UINT16): Likewise.
+	(BT_FN_UINT16_UINT16_UINT16): Likewise.
+	(BT_FN_INT_PTR_INT): Likewise.
+	(BT_FN_UINT_PTR_UINT): Likewise.
+	(BT_FN_LONG_PTR_LONG): Likewise.
+	(BT_FN_ULONG_PTR_ULONG): Likewise.
+	(BT_FN_VOID_UINT64_UINT64): Likewise.
+	(BT_FN_UINT8_UINT8_UINT8): Likewise.
+	(BT_FN_INT8_INT8_INT8): Likewise.
+	(BT_FN_INT16_INT16_INT16): Likewise.
+	(BT_FN_INT_INT_INT): Likewise.
+	(BT_FN_UINT_FLOAT_UINT): Likewise.
+	(BT_FN_FLOAT_UINT_UINT): Likewise.
+	(BT_FN_ULONG_UINT_UINT): Likewise.
+	(BT_FN_ULONG_UINT_PTR): Likewise.
+	(BT_FN_ULONG_ULONG_ULONG): Likewise.
+	(BT_FN_UINT_UINT_UINT): Likewise.
+	(BT_FN_VOID_UINT_PTR): Likewise.
+	(BT_FN_UINT_UINT_PTR: Likewise.
+	(BT_FN_UINT32_UINT64_PTR): Likewise.
+	(BT_FN_INT_INT_UINT_UINT): Likewise.
+	(BT_FN_UINT_UINT_UINT_UINT): Likewise.
+	(BT_FN_UINT_UINT_UINT_PTR): Likewise.
+	(BT_FN_UINT_ULONG_ULONG_UINT): Likewise.
+	(BT_FN_ULONG_ULONG_ULONG_ULONG): Likewise.
+	(BT_FN_LONG_LONG_UINT_UINT): Likewise.
+	(BT_FN_ULONG_ULONG_UINT_UINT): Likewise.
+	(BT_FN_VOID_UINT32_UINT64_PTR): Likewise.
+	(BT_FN_VOID_UINT32_UINT32_PTR): Likewise.
+	(BT_FN_UINT_UINT_UINT_UINT_UINT): Likewise.
+	(BT_FN_UINT_FLOAT_FLOAT_FLOAT_FLOAT): Likewise.
+	(BT_FN_ULONG_ULONG_ULONG_UINT_UINT): Likewise.
+	* doc/frontends.texi: List BRIG FE.
+	* doc/install.texi (Testing): Add BRIG tesring requirements.
+	* doc/invoke.texi (Overall Options): Mention BRIG.
+	* doc/standards.texi (Standards): Doucment BRIG HSA version.
+
 2017-01-24  Richard Biener  <rguenther@suse.de>
 
 	PR translation/79208

gcc/brig/ChangeLog

+2017-01-24  Pekka Jääskeläinen <pekka@parmance.com>
+	    Martin Jambor  <mjambor@suse.cz>
+
+	* Make-lang.in: New file.
+	* brig-builtins.h: Likewise.
+	* brig-c.h: Likewise.
+	* brig-lang.c: Likewise.
+	* brigspec.c: Likewise.
+	* config-lang.in: Likewise.
+	* lang-specs.h: Likewise.
+	* lang.opt: Likewise.
+	* brigfrontend/brig-arg-block-handler.cc: Likewise.
+	* brigfrontend/brig-atomic-inst-handler.cc: Likewise.
+	* brigfrontend/brig-basic-inst-handler.cc: Likewise.
+	* brigfrontend/brig-branch-inst-handler.cc: Likewise.
+	* brigfrontend/brig-cmp-inst-handler.cc: Likewise.
+	* brigfrontend/brig-code-entry-handler.cc: Likewise.
+	* brigfrontend/brig-code-entry-handler.h: Likewise.
+	* brigfrontend/brig-comment-handler.cc: Likewise.
+	* brigfrontend/brig-control-handler.cc: Likewise.
+	* brigfrontend/brig-copy-move-inst-handler.cc: Likewise.
+	* brigfrontend/brig-cvt-inst-handler.cc: Likewise.
+	* brigfrontend/brig-fbarrier-handler.cc: Likewise.
+	* brigfrontend/brig-function-handler.cc: Likewise.
+	* brigfrontend/brig-function.cc: Likewise.
+	* brigfrontend/brig-function.h: Likewise.
+	* brigfrontend/brig-inst-mod-handler.cc: Likewise.
+	* brigfrontend/brig-label-handler.cc: Likewise.
+	* brigfrontend/brig-lane-inst-handler.cc: Likewise.
+	* brigfrontend/brig-machine.c: Likewise.
+	* brigfrontend/brig-machine.h: Likewise.
+	* brigfrontend/brig-mem-inst-handler.cc: Likewise.
+	* brigfrontend/brig-module-handler.cc: Likewise.
+	* brigfrontend/brig-queue-inst-handler.cc: Likewise.
+	* brigfrontend/brig-seg-inst-handler.cc: Likewise.
+	* brigfrontend/brig-signal-inst-handler.cc: Likewise.
+	* brigfrontend/brig-to-generic.cc: Likewise.
+	* brigfrontend/brig-to-generic.h: Likewise.
+	* brigfrontend/brig-util.cc: Likewise.
+	* brigfrontend/brig-util.h: Likewise.
+	* brigfrontend/brig-variable-handler.cc: Likewise.
+	* brigfrontend/phsa.h: Likewise.

gcc/brig/Make-lang.in

+# Make-lang.in -- Top level -*- makefile -*- fragment for gcc BRIG (HSAIL)
+# frontend.
+
+# Copyright (C) 2015 Free Software Foundation, Inc.
+
+# This file is part of GCC.
+
+# GCC is free software; you can redistribute it and/or modify
+# it under the terms of the GNU General Public License as published by
+# the Free Software Foundation; either version 3, or (at your option)
+# any later version.
+
+# GCC is distributed in the hope that it will be useful,
+# but WITHOUT ANY WARRANTY; without even the implied warranty of
+# MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+# GNU General Public License for more details.
+
+# You should have received a copy of the GNU General Public License
+# along with GCC; see the file COPYING3.  If not see
+# <http://www.gnu.org/licenses/>.
+
+# This file provides the language dependent support in the main Makefile.
+
+# Installation name.
+
+GCCBRIG_INSTALL_NAME := $(shell echo gccbrig|sed '$(program_transform_name)')
+GCCBRIG_TARGET_INSTALL_NAME := $(target_noncanonical)-$(shell echo gccbrig|sed \
+	'$(program_transform_name)')
+
+# The name for selecting brig in LANGUAGES.
+brig: brig1$(exeext)
+
+.PHONY: brig
+
+CFLAGS-brig/brigspec.o += $(DRIVER_DEFINES)
+
+GCCBRIG_OBJS = $(GCC_OBJS) brig/brigspec.o
+gccbrig$(exeext): $(GCCBRIG_OBJS) $(EXTRA_GCC_OBJS) libcommon-target.a \
+	$(LIBDEPS)
+	+$(LINKER) $(ALL_LINKERFLAGS) $(LDFLAGS) -o $@ \
+	  $(GCCBRIG_OBJS) $(EXTRA_GCC_OBJS) libcommon-target.a \
+	  $(EXTRA_GCC_LIBS) $(LIBS)
+
+# The cross-compiler version.  This is built mainly as a signal to the
+# brig.install-common target.  If this executable exists, it means that
+# brig.all.cross was run.
+gccbrig-cross$(exeext): gccbrig$(exeext)
+	-rm -f gccbrig-cross$(exeext)
+	cp gccbrig$(exeext) gccbrig-cross$(exeext)
+
+# Use strict warnings.
+brig-warn = $(STRICT_WARN)
+
+BRIG_OBJS = \
+	brig/brig-lang.o \
+	brig/brig-code-entry-handler.o \
+	brig/brig-function-handler.o \
+	brig/brig-variable-handler.o \
+	brig/brig-fbarrier-handler.o \
+	brig/brig-label-handler.o \
+	brig/brig-comment-handler.o \
+	brig/brig-basic-inst-handler.o \
+	brig/brig-cvt-inst-handler.o \
+	brig/brig-seg-inst-handler.o \
+	brig/brig-lane-inst-handler.o \
+	brig/brig-queue-inst-handler.o \
+	brig/brig-copy-move-inst-handler.o \
+	brig/brig-signal-inst-handler.o \
+	brig/brig-atomic-inst-handler.o \
+	brig/brig-arg-block-handler.o \
+	brig/brig-control-handler.o \
+	brig/brig-cmp-inst-handler.o \
+	brig/brig-branch-inst-handler.o \
+	brig/brig-mem-inst-handler.o \
+	brig/brig-module-handler.o \
+	brig/brig-inst-mod-handler.o \
+	brig/brig-function.o \
+	brig/brig-to-generic.o \
+	brig/brig-machine.o \
+	brig/brig-util.o
+
+brig_OBJS = $(BRIG_OBJS) brig/brigspec.o
+
+# brig1$(exeext): $(BRIG_OBJS) attribs.o $(BACKEND) $(LIBDEPS)
+# 	+$(LLINKER) $(ALL_LINKERFLAGS) $(LDFLAGS) -o $@ \
+# 	      $(BRIG_OBJS) attribs.o $(BACKEND) $(LIBS) $(BACKENDLIBS)
+
+
+brig1$(exeext): $(BRIG_OBJS) attribs.o $(BACKEND) $(LIBDEPS)
+	+$(LLINKER) $(ALL_LINKERFLAGS) $(LDFLAGS) -o $@ \
+	      $(BRIG_OBJS) attribs.o $(BACKEND) $(LIBS) \
+		  $(BACKENDLIBS)
+
+# Documentation.
+
+GO_TEXI_FILES = \
+	brig/gccbrig.texi \
+	$(gcc_docdir)/include/fdl.texi \
+	$(gcc_docdir)/include/gpl_v3.texi \
+	$(gcc_docdir)/include/gcc-common.texi \
+	gcc-vers.texi
+
+# doc/gccbrig.info: $(BRIG_TEXI_FILES)
+# 	if test "x$(BUILD_INFO)" = xinfo; then \
+# 	  rm -f doc/gccbrig.info*; \
+# 	  $(MAKEINFO) $(MAKEINFOFLAGS) -I $(gcc_docdir) \
+# 		-I $(gcc_docdir)/include -o $@ $<; \
+# 	else true; fi
+
+# doc/gccbrig.dvi: $(BRIG_TEXI_FILES)
+# 	$(TEXI2DVI) -I $(abs_docdir) -I $(abs_docdir)/include -o $@ $<
+
+# doc/gccbrig.pdf: $(BRIG_TEXI_FILES)
+# 	$(TEXI2PDF) -I $(abs_docdir) -I $(abs_docdir)/include -o $@ $<
+
+$(build_htmldir)/brig/index.html: $(BRIG_TEXI_FILES)
+	$(mkinstalldirs) $(@D)
+	rm -f $(@D)/*
+	$(TEXI2HTML) -I $(gcc_docdir) -I $(gcc_docdir)/include \
+		-I $(srcdir)/brig -o $(@D) $<
+
+# Build hooks.
+
+brig.all.cross: gccbrig-cross$(exeext)
+brig.start.encap: gccbrig$(exeext)
+brig.rest.encap:
+#brig.info: doc/gccbrig.info
+brig.info:
+brig.dvi: doc/gccbrig.dvi
+brig.pdf: doc/gccbrig.pdf
+brig.html: $(build_htmldir)/brig/index.html
+brig.srcinfo: #doc/gccbrig.info
+#	-cp -p $^ $(srcdir)/doc
+
+brig.srcextra:
+brig.tags: force
+	cd $(srcdir)/brig; \
+	etags -o TAGS.sub *.c *.h; \
+	etags --include TAGS.sub --include ../TAGS.sub
+brig.man:
+
+#brig.srcman: doc/gccbrig.1
+#	-cp -p $^ $(srcdir)/doc
+
+lang_checks += check-brig
+
+# Install hooks.
+
+brig.install-common: installdirs
+	-rm -f $(DESTDIR)$(bindir)/$(GCCBRIG_INSTALL_NAME)$(exeext)
+	$(INSTALL_PROGRAM) gccbrig$(exeext) \
+	$(DESTDIR)$(bindir)/$(GCCBRIG_INSTALL_NAME)$(exeext)
+	-if test -f brig1$(exeext); then \
+	  if test -f gccbrig-cross$(exeext); then \
+	    :; \
+	  else \
+	    rm -f $(DESTDIR)$(bindir)/$(GCCBRIG_TARGET_INSTALL_NAME)$(exeext); \
+	    ( cd $(DESTDIR)$(bindir) && \
+	      $(LN) $(GCCBRIG_INSTALL_NAME)$(exeext) \
+	      $(GCCBRIG_TARGET_INSTALL_NAME)$(exeext) ); \
+	  fi; \
+	fi
+
+brig.install-plugin:
+
+brig.install-info: #$(DESTDIR)$(infodir)/gccbrig.info
+
+brig.install-pdf: doc/gccbrig.pdf
+	@$(NORMAL_INSTALL)
+	test -z "$(pdfdir)" || $(mkinstalldirs) "$(DESTDIR)$(pdfdir)/gcc"
+	@for p in doc/gccbrig.pdf; do \
+	  if test -f "$$p"; then d=; else d="$(srcdir)/"; fi; \
+	  f=$(pdf__strip_dir) \
+	  echo " $(INSTALL_DATA) '$$d$$p' '$(DESTDIR)$(pdfdir)/gcc/$$f'"; \
+	  $(INSTALL_DATA) "$$d$$p" "$(DESTDIR)$(pdfdir)/gcc/$$f"; \
+	done
+
+brig.install-html: $(build_htmldir)/brig
+	@$(NORMAL_INSTALL)
+	test -z "$(htmldir)" || $(mkinstalldirs) "$(DESTDIR)$(htmldir)"
+	@for p in $(build_htmldir)/brig; do \
+	  if test -f "$$p" || test -d "$$p"; then d=""; else d="$(srcdir)/"; \
+	  fi; \
+	  f=$(html__strip_dir) \
+	  if test -d "$$d$$p"; then \
+	    echo " $(mkinstalldirs) '$(DESTDIR)$(htmldir)/$$f'"; \
+	    $(mkinstalldirs) "$(DESTDIR)$(htmldir)/$$f" || exit 1; \
+	    echo " $(INSTALL_DATA) '$$d$$p'/* '$(DESTDIR)$(htmldir)/$$f'"; \
+	    $(INSTALL_DATA) "$$d$$p"/* "$(DESTDIR)$(htmldir)/$$f"; \
+	  else \
+	    echo " $(INSTALL_DATA) '$$d$$p' '$(DESTDIR)$(htmldir)/$$f'"; \
+	    $(INSTALL_DATA) "$$d$$p" "$(DESTDIR)$(htmldir)/$$f"; \
+	  fi; \
+	done
+
+brig.install-man: #$(DESTDIR)$(man1dir)/$(GCCBRIG_INSTALL_NAME)$(man1ext)
+
+#$(DESTDIR)$(man1dir)/$(GCCBRIG_INSTALL_NAME)$(man1ext): doc/gccbrig.1 \
+#	installdirs
+#	-rm -f $@
+#	-$(INSTALL_DATA) $< $@
+#	-chmod a-x $@
+
+brig.uninstall:
+	rm -rf $(DESTDIR)$(bindir)/$(GCCBRIG_INSTALL_NAME)$(exeext)
+	rm -rf $(DESTDIR)$(man1dir)/$(GCCBRIG_INSTALL_NAME)$(man1ext)
+	rm -rf $(DESTDIR)$(bindir)/$(GCCBRIG_TARGET_INSTALL_NAME)$(exeext)
+	rm -rf $(DESTDIR)$(infodir)/gccbrig.info*
+
+# Clean hooks.
+
+brig.mostlyclean:
+	-rm -f brig/*$(objext)
+	-rm -f brig/*$(coverageexts)
+brig.clean:
+brig.distclean:
+brig.maintainer-clean:
+	-rm -f $(docobjdir)/gccbrig.1
+
+# Stage hooks.
+
+brig.stage1: stage1-start
+	-mv brig/*$(objext) stage1/brig
+brig.stage2: stage2-start
+	-mv brig/*$(objext) stage2/brig
+brig.stage3: stage3-start
+	-mv brig/*$(objext) stage3/brig
+brig.stage4: stage4-start
+	-mv brig/*$(objext) stage4/brig
+brig.stageprofile: stageprofile-start
+	-mv brig/*$(objext) stageprofile/brig
+brig.stagefeedback: stagefeedback-start
+	-mv brig/*$(objext) stagefeedback/brig
+
+CFLAGS-brig/brig-lang.o += -DDEFAULT_TARGET_VERSION=\"$(version)\" \
+	-DDEFAULT_TARGET_MACHINE=\"$(target_noncanonical)\"
+
+BRIGINCLUDES = -I $(srcdir)/brig -I ${HOME}/local/include \
+	-I $(srcdir)/brig/brigfrontend
+
+brig/brig-machine.o: brig/brigfrontend/brig-machine.c
+	$(COMPILE) $(BRIGINCLUDES) $<
+	$(POSTCOMPILE)
+
+brig/%.o: brig/brigfrontend/%.cc
+	$(COMPILE) $(BRIGINCLUDES) $<
+	$(POSTCOMPILE)

gcc/brig/brig-builtins.h

+/* brig-builtins.h -- brig builtin definitions
+   Copyright (C) 2016 Free Software Foundation, Inc.
+
+   Contributed by Pekka Jaaskelainen <pekka.jaaskelainen@parmance.com>
+   for General Processor Tech.
+   This file is part of GCC.
+
+   GCC is free software; you can redistribute it and/or modify it under
+   the terms of the GNU General Public License as published by the Free
+   Software Foundation; either version 3, or (at your option) any later
+   version.
+
+   GCC is distributed in the hope that it will be useful, but WITHOUT ANY
+   WARRANTY; without even the implied warranty of MERCHANTABILITY or
+   FITNESS FOR A PARTICULAR PURPOSE.  See the GNU General Public License
+   for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with GCC; see the file COPYING3.  If not see
+   <http://www.gnu.org/licenses/>.  */
+
+enum built_in_attribute
+{
+#define DEF_ATTR_NULL_TREE(ENUM) ENUM,
+#define DEF_ATTR_INT(ENUM, VALUE) ENUM,
+#define DEF_ATTR_STRING(ENUM, VALUE) ENUM,
+#define DEF_ATTR_IDENT(ENUM, STRING) ENUM,
+#define DEF_ATTR_TREE_LIST(ENUM, PURPOSE, VALUE, CHAIN) ENUM,
+#include "builtin-attrs.def"
+#undef DEF_ATTR_NULL_TREE
+#undef DEF_ATTR_INT
+#undef DEF_ATTR_STRING
+#undef DEF_ATTR_IDENT
+#undef DEF_ATTR_TREE_LIST
+  ATTR_LAST
+};
+
+/* Builtin types.  */
+
+enum brig_builtin_type
+{
+#define DEF_PRIMITIVE_TYPE(NAME, VALUE) NAME,
+#define DEF_FUNCTION_TYPE_0(NAME, RETURN) NAME,
+#define DEF_FUNCTION_TYPE_1(NAME, RETURN, ARG1) NAME,
+#define DEF_FUNCTION_TYPE_2(NAME, RETURN, ARG1, ARG2) NAME,
+#define DEF_FUNCTION_TYPE_3(NAME, RETURN, ARG1, ARG2, ARG3) NAME,
+#define DEF_FUNCTION_TYPE_4(NAME, RETURN, ARG1, ARG2, ARG3, ARG4) NAME,
+#define DEF_FUNCTION_TYPE_5(NAME, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5) NAME,
+#define DEF_FUNCTION_TYPE_6(NAME, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5, \
+			    ARG6) NAME,
+#define DEF_FUNCTION_TYPE_7(NAME, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5, \
+			    ARG6, ARG7) NAME,
+#define DEF_FUNCTION_TYPE_8(NAME, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5, \
+			    ARG6, ARG7, ARG8) NAME,
+#define DEF_FUNCTION_TYPE_9(NAME, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5, \
+			    ARG6, ARG7, ARG8, ARG9) NAME,
+#define DEF_FUNCTION_TYPE_10(NAME, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5, \
+			     ARG6, ARG7, ARG8, ARG9, ARG10) NAME,
+#define DEF_FUNCTION_TYPE_11(NAME, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5, \
+			     ARG6, ARG7, ARG8, ARG9, ARG10, ARG11) NAME,
+#define DEF_FUNCTION_TYPE_VAR_0(NAME, RETURN) NAME,
+#define DEF_FUNCTION_TYPE_VAR_1(NAME, RETURN, ARG1) NAME,
+#define DEF_FUNCTION_TYPE_VAR_2(NAME, RETURN, ARG1, ARG2) NAME,
+#define DEF_FUNCTION_TYPE_VAR_3(NAME, RETURN, ARG1, ARG2, ARG3) NAME,
+#define DEF_FUNCTION_TYPE_VAR_4(NAME, RETURN, ARG1, ARG2, ARG3, ARG4) NAME,
+#define DEF_FUNCTION_TYPE_VAR_5(NAME, RETURN, ARG1, ARG2, ARG3, ARG4, ARG6) \
+				NAME,
+#define DEF_FUNCTION_TYPE_VAR_6(NAME, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5, \
+				 ARG6) NAME,
+#define DEF_FUNCTION_TYPE_VAR_7(NAME, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5, \
+				ARG6, ARG7) NAME,
+#define DEF_POINTER_TYPE(NAME, TYPE) NAME,
+#include "builtin-types.def"
+#undef DEF_PRIMITIVE_TYPE
+#undef DEF_FUNCTION_TYPE_0
+#undef DEF_FUNCTION_TYPE_1
+#undef DEF_FUNCTION_TYPE_2
+#undef DEF_FUNCTION_TYPE_3
+#undef DEF_FUNCTION_TYPE_4
+#undef DEF_FUNCTION_TYPE_5
+#undef DEF_FUNCTION_TYPE_6
+#undef DEF_FUNCTION_TYPE_7
+#undef DEF_FUNCTION_TYPE_8
+#undef DEF_FUNCTION_TYPE_9
+#undef DEF_FUNCTION_TYPE_10
+#undef DEF_FUNCTION_TYPE_11
+#undef DEF_FUNCTION_TYPE_VAR_0
+#undef DEF_FUNCTION_TYPE_VAR_1
+#undef DEF_FUNCTION_TYPE_VAR_2
+#undef DEF_FUNCTION_TYPE_VAR_3
+#undef DEF_FUNCTION_TYPE_VAR_4
+#undef DEF_FUNCTION_TYPE_VAR_5
+#undef DEF_FUNCTION_TYPE_VAR_6
+#undef DEF_FUNCTION_TYPE_VAR_7
+#undef DEF_POINTER_TYPE
+  BT_LAST
+};
+
+typedef enum brig_builtin_type builtin_type;

gcc/brig/brig-c.h

+/* brig-c.h -- Header file for brig input's gcc C interface.
+   Copyright (C) 2016 Free Software Foundation, Inc.
+   Contributed by Pekka Jaaskelainen <pekka.jaaskelainen@parmance.com>
+   for General Processor Tech.
+
+This file is part of GCC.
+
+GCC is free software; you can redistribute it and/or modify it under
+the terms of the GNU General Public License as published by the Free
+Software Foundation; either version 3, or (at your option) any later
+version.
+
+GCC is distributed in the hope that it will be useful, but WITHOUT ANY
+WARRANTY; without even the implied warranty of MERCHANTABILITY or
+FITNESS FOR A PARTICULAR PURPOSE.  See the GNU General Public License
+for more details.
+
+You should have received a copy of the GNU General Public License
+along with GCC; see the file COPYING3.  If not see
+<http://www.gnu.org/licenses/>.  */
+
+#ifndef BRIG_BRIG_C_H
+#define BRIG_BRIG_C_H
+
+#define BRIG_EXTERN_C
+
+#include "machmode.h"
+
+/* Functions defined in the Brig frontend proper called by the GCC
+   interface.  */
+
+extern int brig_enable_dump (const char *);
+extern int brig_enable_optimize (const char *);
+
+extern void brig_add_search_path (const char *);
+
+extern void brig_create_brigbrig (int int_type_size, int pointer_size,
+				  const char *pkgpath, const char *prefix,
+				  const char *relative_import_path);
+
+extern void brig_parse_input_files (const char **, unsigned int,
+				    bool only_check_syntax,
+				    bool require_return_statement);
+extern void brig_write_globals (void);
+
+extern tree brig_type_for_size (unsigned int bits, int unsignedp);
+extern tree brig_type_for_mode (enum machine_mode, int unsignedp);
+
+/* Functions defined in the GCC interface called by the Brig frontend
+   proper.  */
+
+extern void brig_preserve_from_gc (tree);
+
+extern const char *brig_localize_identifier (const char *);
+
+extern unsigned int brig_field_alignment (tree);
+
+extern void brig_trampoline_info (unsigned int *size, unsigned int *alignment);
+
+extern void brig_imported_unsafe (void);
+
+extern void brig_write_export_data (const char *, unsigned int);
+
+extern const char *brig_read_export_data (int, off_t, char **, size_t *, int *);
+
+#endif /* !defined (BRIG_BRIG_C_H) */

gcc/brig/brigfrontend/brig-arg-block-handler.cc

+/* brig-arg-block-handler.cc -- brig arg block start/end directive handling
+   Copyright (C) 2016 Free Software Foundation, Inc.
+   Contributed by Pekka Jaaskelainen <pekka.jaaskelainen@parmance.com>
+   for General Processor Tech.
+
+   This file is part of GCC.
+
+   GCC is free software; you can redistribute it and/or modify it under
+   the terms of the GNU General Public License as published by the Free
+   Software Foundation; either version 3, or (at your option) any later
+   version.
+
+   GCC is distributed in the hope that it will be useful, but WITHOUT ANY
+   WARRANTY; without even the implied warranty of MERCHANTABILITY or
+   FITNESS FOR A PARTICULAR PURPOSE.  See the GNU General Public License
+   for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with GCC; see the file COPYING3.  If not see
+   <http://www.gnu.org/licenses/>.  */
+
+#include "brig-code-entry-handler.h"
+#include "tree-iterator.h"
+#include "system.h"
+#include "errors.h"
+
+#include "tree-pretty-print.h"
+#include "print-tree.h"
+
+size_t
+brig_directive_arg_block_handler::operator () (const BrigBase *base)
+{
+  if (base->kind == BRIG_KIND_DIRECTIVE_ARG_BLOCK_START)
+    {
+      /* Initiate a new code block for the call site.  */
+      tree stmt_list = alloc_stmt_list ();
+      tree bind_expr
+	= build3 (BIND_EXPR, void_type_node, NULL, stmt_list, NULL);
+      tree block = make_node (BLOCK);
+      BIND_EXPR_BLOCK (bind_expr) = block;
+      static int block_id = 0;
+      BLOCK_NUMBER (block) = block_id++;
+      TREE_USED (block) = 1;
+      tree m_parentblock = DECL_INITIAL (m_parent.m_cf->m_func_decl);
+      BLOCK_SUPERCONTEXT (block) = m_parentblock;
+
+      chainon (BLOCK_SUBBLOCKS (m_parentblock), block);
+
+      m_parent.m_cf->m_current_bind_expr = bind_expr;
+      m_parent.m_cf->m_generating_arg_block = true;
+    }
+  else if (base->kind == BRIG_KIND_DIRECTIVE_ARG_BLOCK_END)
+    {
+      /* Restore the used bind expression back to the function
+	 scope.  */
+      tree new_bind_expr = m_parent.m_cf->m_current_bind_expr;
+      m_parent.m_cf->m_current_bind_expr
+	= DECL_SAVED_TREE (m_parent.m_cf->m_func_decl);
+      m_parent.m_cf->append_statement (new_bind_expr);
+      m_parent.m_cf->m_generating_arg_block = false;
+    }
+  else
+    gcc_unreachable ();
+
+  return base->byteCount;
+}

gcc/brig/brigfrontend/brig-atomic-inst-handler.cc

+/* brig-atomic-inst-handler.cc -- brig atomic instruction handling
+   Copyright (C) 2016 Free Software Foundation, Inc.
+
+   Contributed by Pekka Jaaskelainen <pekka.jaaskelainen@parmance.com>
+   for General Processor Tech.
+   This file is part of GCC.
+
+   GCC is free software; you can redistribute it and/or modify it under
+   the terms of the GNU General Public License as published by the Free
+   Software Foundation; either version 3, or (at your option) any later
+   version.
+
+   GCC is distributed in the hope that it will be useful, but WITHOUT ANY
+   WARRANTY; without even the implied warranty of MERCHANTABILITY or
+   FITNESS FOR A PARTICULAR PURPOSE.  See the GNU General Public License
+   for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with GCC; see the file COPYING3.  If not see
+   <http://www.gnu.org/licenses/>.  */
+
+#include <sstream>
+
+#include "brig-code-entry-handler.h"
+#include "brig-util.h"
+#include "fold-const.h"
+#include "diagnostic.h"
+#include "tree-pretty-print.h"
+#include "print-tree.h"
+#include "convert.h"
+#include "langhooks.h"
+#include "gimple-expr.h"
+#include "stringpool.h"
+#include "brig-builtins.h"
+
+brig_atomic_inst_handler::brig_atomic_inst_handler (brig_to_generic &parent)
+  : brig_code_entry_handler (parent)
+{
+}
+
+size_t
+brig_atomic_inst_handler::generate_tree (const BrigInstBase &inst,
+					 BrigAtomicOperation8_t atomic_opcode)
+{
+  tree_stl_vec operands = build_operands (inst);
+  const int first_input
+    = gccbrig_hsa_opcode_op_output_p (inst.opcode, 0) ? 1 : 0;
+
+  tree instr_type = gccbrig_tree_type_for_hsa_type (inst.type);
+
+  /* Utilize the atomic data types (from C++11 support) for implementing
+     atomic operations.  */
+
+  tree atomic_type = build_qualified_type (instr_type, TYPE_QUAL_ATOMIC);
+
+  gcc_assert (atomic_type != NULL_TREE);
+
+  tree signal_handle = operands[first_input];
+  tree atomic_ptype = build_pointer_type (atomic_type);
+  tree casted_to_ptr = convert_to_pointer (atomic_ptype, signal_handle);
+
+  tree src0 = NULL_TREE;
+  if (atomic_opcode != BRIG_ATOMIC_LD)
+    src0 = operands[first_input + 1];
+
+  tree instr_expr = NULL_TREE;
+
+  tree ptype = build_pointer_type (instr_type);
+  tree ptr = convert_to_pointer (ptype, operands[first_input]);
+
+  if (atomic_opcode == BRIG_ATOMIC_ST)
+    {
+      tree mem_ref = build2 (MEM_REF, atomic_type, casted_to_ptr,
+			     build_int_cst (atomic_ptype, 0));
+      instr_expr = build2 (MODIFY_EXPR, atomic_type, mem_ref, src0);
+    }
+  else if (atomic_opcode == BRIG_ATOMIC_LD
+	   || (atomic_opcode >= BRIG_ATOMIC_WAIT_EQ
+	       && atomic_opcode <= BRIG_ATOMIC_WAITTIMEOUT_GTE))
+    {
+      tree mem_ref = build2 (MEM_REF, atomic_type, casted_to_ptr,
+			     build_int_cst (atomic_ptype, 0));
+      /* signal_wait* instructions can return spuriously before the
+	 condition becomes true.  Therefore it's legal to return
+	 right away.  TODO: builtin calls which can be
+	 implemented with a power efficient sleep-wait.  */
+      instr_expr = mem_ref;
+    }
+  else if (atomic_opcode == BRIG_ATOMIC_CAS)
+    {
+      /* Special case for CAS due to the two args.  */
+      tree built_in = NULL_TREE;
+      switch (gccbrig_hsa_type_bit_size (inst.type))
+	{
+	case 32:
+	  built_in
+	    = builtin_decl_explicit (BUILT_IN_SYNC_VAL_COMPARE_AND_SWAP_4);
+	  break;
+	case 64:
+	  built_in
+	    = builtin_decl_explicit (BUILT_IN_SYNC_VAL_COMPARE_AND_SWAP_8);
+	  break;
+	default:
+	  gcc_unreachable ();
+	}
+
+      tree src1 = operands[first_input + 2];
+
+      tree src0_type
+	= TREE_VALUE (TREE_CHAIN (TYPE_ARG_TYPES (TREE_TYPE (built_in))));
+
+      tree src1_type = TREE_VALUE
+	(TREE_CHAIN (TREE_CHAIN (TYPE_ARG_TYPES (TREE_TYPE (built_in)))));
+
+      instr_expr = call_builtin (built_in, 3, instr_type, ptype, ptr,
+				 src0_type, src0, src1_type, src1);
+    }
+  else
+    {
+      tree built_in = NULL_TREE;
+      /* The rest of the builtins have the same number of parameters.
+	 Generate a big if..else that finds the correct builtin
+	 automagically from the def file.  */
+#undef DEF_HSAIL_SAT_BUILTIN
+#undef DEF_HSAIL_BUILTIN
+#undef DEF_HSAIL_ATOMIC_BUILTIN
+#undef DEF_HSAIL_INTR_BUILTIN
+#undef DEF_HSAIL_CVT_ZEROI_SAT_BUILTIN
+
+#define DEF_HSAIL_ATOMIC_BUILTIN(ENUM, ATOMIC_OPCODE, HSAIL_TYPE,	\
+				 NAME, TYPE, ATTRS)			\
+      if (atomic_opcode == ATOMIC_OPCODE && inst.type == HSAIL_TYPE)	\
+	built_in = builtin_decl_explicit (ENUM);			\
+      else
+#include "brig-builtins.def"
+      switch (atomic_opcode)
+	{
+	case BRIG_ATOMIC_ADD:
+	  switch (gccbrig_hsa_type_bit_size (inst.type))
+	    {
+	    case 32:
+	      built_in
+		= builtin_decl_explicit (BUILT_IN_SYNC_FETCH_AND_ADD_4);
+	      break;
+	    case 64:
+	      built_in
+		= builtin_decl_explicit (BUILT_IN_SYNC_FETCH_AND_ADD_8);
+	      break;
+	    default:
+	      gcc_unreachable ();
+	    }
+	  break;
+	case BRIG_ATOMIC_SUB:
+	  switch (gccbrig_hsa_type_bit_size (inst.type))
+	    {
+	    case 32:
+	      built_in
+		= builtin_decl_explicit (BUILT_IN_SYNC_FETCH_AND_SUB_4);
+	      break;
+	    case 64:
+	      built_in
+		= builtin_decl_explicit (BUILT_IN_SYNC_FETCH_AND_SUB_8);
+	      break;
+	    default:
+	      gcc_unreachable ();
+	    }
+	  break;
+	case BRIG_ATOMIC_AND:
+	  switch (gccbrig_hsa_type_bit_size (inst.type))
+	    {
+	    case 32:
+	      built_in
+		= builtin_decl_explicit (BUILT_IN_SYNC_FETCH_AND_AND_4);
+	      break;
+	    case 64:
+	      built_in
+		= builtin_decl_explicit (BUILT_IN_SYNC_FETCH_AND_AND_8);
+	      break;
+	    default:
+	      gcc_unreachable ();
+	    }
+	  break;
+	case BRIG_ATOMIC_XOR:
+	  switch (gccbrig_hsa_type_bit_size (inst.type))
+	    {
+	    case 32:
+	      built_in
+		= builtin_decl_explicit (BUILT_IN_SYNC_FETCH_AND_XOR_4);
+	      break;
+	    case 64:
+	      built_in
+		= builtin_decl_explicit (BUILT_IN_SYNC_FETCH_AND_XOR_8);
+	      break;
+	    default:
+	      gcc_unreachable ();
+	    }
+	  break;
+	case BRIG_ATOMIC_OR:
+	  switch (gccbrig_hsa_type_bit_size (inst.type))
+	    {
+	    case 32:
+	      built_in
+		= builtin_decl_explicit (BUILT_IN_SYNC_FETCH_AND_OR_4);
+	      break;
+	    case 64:
+	      built_in
+		= builtin_decl_explicit (BUILT_IN_SYNC_FETCH_AND_OR_8);
+	      break;
+	    default:
+	      gcc_unreachable ();
+	    }
+	  break;
+	case BRIG_ATOMIC_EXCH:
+	  switch (gccbrig_hsa_type_bit_size (inst.type))
+	    {
+	    case 32:
+	      built_in
+		= builtin_decl_explicit (BUILT_IN_SYNC_LOCK_TEST_AND_SET_4);
+	      break;
+	    case 64:
+	      built_in
+		= builtin_decl_explicit (BUILT_IN_SYNC_LOCK_TEST_AND_SET_8);
+	      break;
+	    default:
+	      gcc_unreachable ();
+	    }
+	  break;
+	default:
+	  gcc_unreachable ();
+	};
+
+      gcc_assert (built_in != NULL_TREE);
+      tree arg0_type
+	= TREE_VALUE (TREE_CHAIN (TYPE_ARG_TYPES (TREE_TYPE (built_in))));
+
+      instr_expr = call_builtin (built_in, 2, instr_type, ptr_type_node,
+				 ptr, arg0_type, src0);
+
+      /* We need a temp variable for the result, because otherwise
+	 the gimplifier drops a necessary (unsigned to signed) cast in
+	 the output assignment and fails a check later.  */
+      tree tmp_var = create_tmp_var (arg0_type, "builtin_out");
+      tree tmp_assign
+	= build2 (MODIFY_EXPR, TREE_TYPE (tmp_var), tmp_var, instr_expr);
+      m_parent.m_cf->append_statement (tmp_assign);
+      instr_expr = tmp_var;
+    }
+
+  if (first_input > 0)
+    build_output_assignment (inst, operands[0], instr_expr);
+  else
+    m_parent.m_cf->append_statement (instr_expr);
+
+  return inst.base.byteCount;
+}
+
+size_t
+brig_atomic_inst_handler::operator () (const BrigBase *base)
+{
+  const BrigInstAtomic *inst = (const BrigInstAtomic *) base;
+  BrigAtomicOperation8_t atomic_opcode;
+  atomic_opcode = inst->atomicOperation;
+
+  return generate_tree (inst->base, atomic_opcode);
+}

gcc/brig/brigfrontend/brig-branch-inst-handler.cc

+/* brig-branch-inst-handler.cc -- brig branch instruction handling
+   Copyright (C) 2016 Free Software Foundation, Inc.
+   Contributed by Pekka Jaaskelainen <pekka.jaaskelainen@parmance.com>
+   for General Processor Tech.
+
+   This file is part of GCC.
+
+   GCC is free software; you can redistribute it and/or modify it under
+   the terms of the GNU General Public License as published by the Free
+   Software Foundation; either version 3, or (at your option) any later
+   version.
+
+   GCC is distributed in the hope that it will be useful, but WITHOUT ANY
+   WARRANTY; without even the implied warranty of MERCHANTABILITY or
+   FITNESS FOR A PARTICULAR PURPOSE.  See the GNU General Public License
+   for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with GCC; see the file COPYING3.  If not see
+   <http://www.gnu.org/licenses/>.  */
+
+#include "brig-code-entry-handler.h"
+
+#include "errors.h"
+#include "brig-util.h"
+#include "tree-pretty-print.h"
+#include "print-tree.h"
+#include "vec.h"
+#include "fold-const.h"
+
+size_t
+brig_branch_inst_handler::operator () (const BrigBase *base)
+{
+  const BrigInstBase *brig_inst
+    = (const BrigInstBase *) &((const BrigInstBasic *) base)->base;
+
+  if (brig_inst->opcode == BRIG_OPCODE_CALL)
+    {
+      const BrigData *operand_entries
+	= m_parent.get_brig_data_entry (brig_inst->operands);
+      tree func_ref = NULL_TREE;
+      vec<tree, va_gc> *out_args;
+      vec_alloc (out_args, 1);
+      vec<tree, va_gc> *in_args;
+      vec_alloc (in_args, 4);
+
+      size_t operand_count = operand_entries->byteCount / 4;
+      gcc_assert (operand_count < 4);
+
+      for (size_t i = 0; i < operand_count; ++i)
+	{
+	  uint32_t operand_offset
+	    = ((const uint32_t *) &operand_entries->bytes)[i];
+	  const BrigBase *operand_data
+	    = m_parent.get_brig_operand_entry (operand_offset);
+	  if (i == 1)
+	    {
+	      gcc_assert (operand_data->kind == BRIG_KIND_OPERAND_CODE_REF);
+	      func_ref = build_tree_operand (*brig_inst, *operand_data);
+	      continue;
+	    }
+	  gcc_assert (operand_data->kind == BRIG_KIND_OPERAND_CODE_LIST);
+	  const BrigOperandCodeList *codelist
+	    = (const BrigOperandCodeList *) operand_data;
+	  const BrigData *data
+	    = m_parent.get_brig_data_entry (codelist->elements);
+
+	  size_t bytes = data->byteCount;
+	  const BrigOperandOffset32_t *operand_ptr
+	    = (const BrigOperandOffset32_t *) data->bytes;
+
+	  vec<tree, va_gc> *args = i == 0 ? out_args : in_args;
+
+	  while (bytes > 0)
+	    {
+	      BrigOperandOffset32_t offset = *operand_ptr;
+	      const BrigBase *code_element
+		= m_parent.get_brig_code_entry (offset);
+	      gcc_assert (code_element->kind == BRIG_KIND_DIRECTIVE_VARIABLE);
+	      const BrigDirectiveVariable *brig_var
+		= (const BrigDirectiveVariable *) code_element;
+	      tree var = m_parent.m_cf->arg_variable (brig_var);
+
+	      if (brig_var->type & BRIG_TYPE_ARRAY)
+		{
+		  /* Array return values are passed as the first argument.  */
+		  args = in_args;
+		  /* Pass pointer to the element zero and use its element zero
+		     as the base address.  */
+		  tree etype = TREE_TYPE (TREE_TYPE (var));
+		  tree ptype = build_pointer_type (etype);
+		  tree element_zero
+		    = build4 (ARRAY_REF, etype, var, integer_zero_node,
+			      NULL_TREE, NULL_TREE);
+		  var = build1 (ADDR_EXPR, ptype, element_zero);
+		}
+
+	      gcc_assert (var != NULL_TREE);
+	      vec_safe_push (args, var);
+	      ++operand_ptr;
+	      bytes -= 4;
+	    }
+	}
+
+      gcc_assert (func_ref != NULL_TREE);
+      gcc_assert (out_args->length () == 0 || out_args->length () == 1);
+
+      tree ret_val_type = void_type_node;
+      tree ret_val = NULL_TREE;
+      if (out_args->length () == 1)
+	{
+	  ret_val = (*out_args)[0];
+	  ret_val_type = TREE_TYPE (ret_val);
+	}
+
+      /* Pass the hidden kernel arguments along to the called functions as
+	 they might call builtins that need them or access group/private
+	 memory.  */
+
+      vec_safe_push (in_args, m_parent.m_cf->m_context_arg);
+      vec_safe_push (in_args, m_parent.m_cf->m_group_base_arg);
+      vec_safe_push (in_args, m_parent.m_cf->m_private_base_arg);
+
+      tree call = build_call_vec (ret_val_type, build_fold_addr_expr (func_ref),
+				  in_args);
+      TREE_NOTHROW (func_ref) = 1;
+      TREE_NOTHROW (call) = 1;
+
+      if (ret_val != NULL_TREE)
+	{
+	  TREE_ADDRESSABLE (ret_val) = 1;
+	  tree result_assign
+	    = build2 (MODIFY_EXPR, TREE_TYPE (ret_val), ret_val, call);
+	  m_parent.m_cf->append_statement (result_assign);
+	}
+      else
+	{
+	  m_parent.m_cf->append_statement (call);
+	}
+
+      m_parent.m_cf->m_has_unexpanded_dp_builtins = false;
+      m_parent.m_cf->m_called_functions.push_back (func_ref);
+
+      return base->byteCount;
+    }
+
+  tree instr_type = gccbrig_tree_type_for_hsa_type (brig_inst->type);
+  tree_stl_vec operands = build_operands (*brig_inst);
+
+  if (brig_inst->opcode == BRIG_OPCODE_BR)
+    {
+      tree goto_stmt = build1 (GOTO_EXPR, instr_type, operands[0]);
+      m_parent.m_cf->append_statement (goto_stmt);
+    }
+  else if (brig_inst->opcode == BRIG_OPCODE_SBR)
+    {
+      tree select = operands[0];
+      tree cases = operands[1];
+
+      tree switch_expr = build3 (SWITCH_EXPR, TREE_TYPE (select), select,
+				 NULL_TREE, NULL_TREE);
+
+      tree default_case
+	= build_case_label (NULL_TREE, NULL_TREE,
+			    create_artificial_label (UNKNOWN_LOCATION));
+      append_to_statement_list (default_case, &SWITCH_BODY (switch_expr));
+
+      tree default_jump
+	= build1 (GOTO_EXPR, void_type_node, TREE_VEC_ELT (cases, 0));
+      append_to_statement_list (default_jump, &SWITCH_BODY (switch_expr));
+
+      for (int c = 0; c < TREE_VEC_LENGTH (cases); ++c)
+	{
+	  tree case_label
+	    = build_case_label (build_int_cst (integer_type_node, c), NULL_TREE,
+				create_artificial_label (UNKNOWN_LOCATION));
+
+	  append_to_statement_list (case_label, &SWITCH_BODY (switch_expr));
+
+	  tree jump
+	    = build1 (GOTO_EXPR, void_type_node, TREE_VEC_ELT (cases, c));
+	  append_to_statement_list (jump, &SWITCH_BODY (switch_expr));
+	}
+      m_parent.m_cf->append_statement (switch_expr);
+    }
+  else if (brig_inst->opcode == BRIG_OPCODE_CBR)
+    {
+      tree condition = operands[0];
+      tree target_goto = build1 (GOTO_EXPR, void_type_node, operands[1]);
+      /* Represents the if..else as (condition)?(goto foo):(goto bar).  */
+      tree if_stmt
+	= build3 (COND_EXPR, void_type_node, condition, target_goto, NULL_TREE);
+      m_parent.m_cf->append_statement (if_stmt);
+    }
+  else if (brig_inst->opcode == BRIG_OPCODE_WAVEBARRIER)
+    {
+      /* WAVEBARRIER is a NOP when WAVESIZE = 1.  */
+    }
+  else if (brig_inst->opcode == BRIG_OPCODE_BARRIER)
+    {
+      m_parent.m_cf->m_has_barriers = true;
+      tree_stl_vec call_operands;
+      /* FIXME.  We should add attributes (are there suitable ones in gcc?) that
+	 ensure the barrier won't be duplicated or moved out of loops etc.
+	 Like the 'noduplicate' of LLVM.  Same goes for fbarriers.  */
+      m_parent.m_cf->append_statement
+	(expand_or_call_builtin (brig_inst->opcode, BRIG_TYPE_NONE, NULL_TREE,
+				 call_operands));
+    }
+  else if (brig_inst->opcode >= BRIG_OPCODE_ARRIVEFBAR
+	   && brig_inst->opcode <= BRIG_OPCODE_WAITFBAR)
+    {
+      m_parent.m_cf->m_has_barriers = true;
+      m_parent.m_cf->append_statement
+	(expand_or_call_builtin (brig_inst->opcode, BRIG_TYPE_NONE,
+				 uint32_type_node, operands));
+    }
+  else
+    gcc_unreachable ();
+  return base->byteCount;
+}

gcc/brig/brigfrontend/brig-cmp-inst-handler.cc

+/* brig-cmp-inst-handler.cc -- brig cmp instruction handling
+   Copyright (C) 2016 Free Software Foundation, Inc.
+   Contributed by Pekka Jaaskelainen <pekka.jaaskelainen@parmance.com>
+   for General Processor Tech.
+
+This file is part of GCC.
+
+GCC is free software; you can redistribute it and/or modify it under
+the terms of the GNU General Public License as published by the Free
+Software Foundation; either version 3, or (at your option) any later
+version.
+
+GCC is distributed in the hope that it will be useful, but WITHOUT ANY
+WARRANTY; without even the implied warranty of MERCHANTABILITY or
+FITNESS FOR A PARTICULAR PURPOSE.  See the GNU General Public License
+for more details.
+
+You should have received a copy of the GNU General Public License
+along with GCC; see the file COPYING3.  If not see
+<http://www.gnu.org/licenses/>.  */
+
+#include "brig-code-entry-handler.h"
+#include "diagnostic.h"
+#include "tree-pretty-print.h"
+#include "print-tree.h"
+#include "brig-util.h"
+#include "convert.h"
+
+size_t
+brig_cmp_inst_handler::operator () (const BrigBase *base)
+{
+  const BrigInstBase *inst_base = (const BrigInstBase *) base;
+  const BrigInstCmp *inst = (const BrigInstCmp *) base;
+
+  tree cmp_type = get_tree_expr_type_for_hsa_type (inst->sourceType);
+
+  /* The destination type to convert the comparison result to.  */
+  tree dest_type = gccbrig_tree_type_for_hsa_type (inst_base->type);
+
+  const bool is_fp16_dest
+    = (inst_base->type & BRIG_TYPE_BASE_MASK) == BRIG_TYPE_F16;
+  const bool is_boolean_dest
+    = (inst_base->type & BRIG_TYPE_BASE_MASK) == BRIG_TYPE_B1;
+
+  bool is_int_cmp = VECTOR_TYPE_P (cmp_type)
+		      ? INTEGRAL_TYPE_P (TREE_TYPE (cmp_type))
+		      : INTEGRAL_TYPE_P (cmp_type);
+
+  /* The type for the GENERIC comparison.  It should match the
+     input operand width for vector comparisons, a boolean
+     otherwise.  */
+  tree result_type = get_comparison_result_type (cmp_type);
+
+  /* Save the result as a boolean and extend/convert it to the
+     wanted destination type.  */
+  tree expr = NULL_TREE;
+
+  std::vector<tree> operands = build_operands (*inst_base);
+
+  switch (inst->compare)
+    {
+    case BRIG_COMPARE_SEQ:
+    case BRIG_COMPARE_EQ:
+      expr = build2 (EQ_EXPR, result_type, operands[1], operands[2]);
+      break;
+    case BRIG_COMPARE_SNE:
+    case BRIG_COMPARE_NE:
+      expr = build2 (NE_EXPR, result_type, operands[1], operands[2]);
+
+      if (!is_int_cmp)
+	expr = build2 (BIT_AND_EXPR, TREE_TYPE (expr),
+		       expr,
+		       build2 (ORDERED_EXPR, result_type, operands[1],
+			       operands[2]));
+      break;
+    case BRIG_COMPARE_SLT:
+    case BRIG_COMPARE_LT:
+      expr = build2 (LT_EXPR, result_type, operands[1], operands[2]);
+      break;
+    case BRIG_COMPARE_SLE:
+    case BRIG_COMPARE_LE:
+      expr = build2 (LE_EXPR, result_type, operands[1], operands[2]);
+      break;
+    case BRIG_COMPARE_SGT:
+    case BRIG_COMPARE_GT:
+      expr = build2 (GT_EXPR, result_type, operands[1], operands[2]);
+      break;
+    case BRIG_COMPARE_SGE:
+    case BRIG_COMPARE_GE:
+      expr = build2 (GE_EXPR, result_type, operands[1], operands[2]);
+      break;
+    case BRIG_COMPARE_SEQU:
+    case BRIG_COMPARE_EQU:
+      expr = build2 (UNEQ_EXPR, result_type, operands[1], operands[2]);
+      break;
+    case BRIG_COMPARE_SNEU:
+    case BRIG_COMPARE_NEU:
+      expr = build2 (NE_EXPR, result_type, operands[1], operands[2]);
+      break;
+    case BRIG_COMPARE_SLTU:
+    case BRIG_COMPARE_LTU:
+      expr = build2 (UNLT_EXPR, result_type, operands[1], operands[2]);
+      break;
+    case BRIG_COMPARE_SLEU:
+    case BRIG_COMPARE_LEU:
+      expr = build2 (UNLE_EXPR, result_type, operands[1], operands[2]);
+      break;
+    case BRIG_COMPARE_SGTU:
+    case BRIG_COMPARE_GTU:
+      expr = build2 (UNGT_EXPR, result_type, operands[1], operands[2]);
+      break;
+    case BRIG_COMPARE_SGEU:
+    case BRIG_COMPARE_GEU:
+      expr = build2 (UNGE_EXPR, result_type, operands[1], operands[2]);
+      break;
+    case BRIG_COMPARE_SNUM:
+    case BRIG_COMPARE_NUM:
+      expr = build2 (ORDERED_EXPR, result_type, operands[1], operands[2]);
+      break;
+    case BRIG_COMPARE_SNAN:
+    case BRIG_COMPARE_NAN:
+      expr = build2 (UNORDERED_EXPR, result_type, operands[1], operands[2]);
+      break;
+    default:
+      break;
+    }
+
+  if (expr == NULL_TREE)
+    gcc_unreachable ();
+
+  if (is_fp16_dest)
+    {
+      expr = convert_to_real (brig_to_generic::s_fp32_type, expr);
+    }
+  else if (VECTOR_TYPE_P (dest_type) && ANY_INTEGRAL_TYPE_P (dest_type)
+	   && !is_boolean_dest
+	   && (inst->sourceType & BRIG_TYPE_BASE_MASK) != BRIG_TYPE_F16)
+    {
+      /* In later gcc versions, the output of comparison is not
+	 all ones for vectors like still in 4.9.1.  We need to use
+	 an additional VEC_COND_EXPR to produce the all ones 'true' value
+	 required by HSA.
+	 VEC_COND_EXPR <a == b, { -1, -1, -1, -1 }, { 0, 0, 0, 0 }>; */
+
+      tree all_ones
+	= build_vector_from_val (dest_type,
+			       build_minus_one_cst (TREE_TYPE (dest_type)));
+      tree all_zeroes
+	= build_vector_from_val (dest_type,
+			       build_zero_cst (TREE_TYPE (dest_type)));
+      expr = build3 (VEC_COND_EXPR, dest_type, expr, all_ones, all_zeroes);
+    }
+  else if (INTEGRAL_TYPE_P (dest_type) && !is_boolean_dest)
+    {
+      /* We need to produce the all-ones pattern for the width of the whole
+	 resulting integer type.  Use back and forth shifts for propagating
+	 the lower 1.  */
+      tree signed_type = signed_type_for (dest_type);
+      tree signed_result = convert_to_integer (signed_type, expr);
+
+      size_t result_width = int_size_in_bytes (dest_type) * BITS_PER_UNIT;
+
+      tree shift_amount_cst
+	= build_int_cstu (signed_type, result_width - 1);
+
+      tree shift_left_result
+	= build2 (LSHIFT_EXPR, signed_type, signed_result, shift_amount_cst);
+
+      expr = build2 (RSHIFT_EXPR, signed_type, shift_left_result,
+		     shift_amount_cst);
+    }
+  else if (SCALAR_FLOAT_TYPE_P (dest_type))
+    {
+      expr = convert_to_real (dest_type, expr);
+    }
+  else if (VECTOR_TYPE_P (dest_type)
+	   && (inst->sourceType & BRIG_TYPE_BASE_MASK) == BRIG_TYPE_F16)
+    {
+      /* Because F16 comparison is emulated as an F32 comparison with S32
+	 results, we must now truncate the result vector to S16s so it
+	 fits to the destination register.  We can build the target vector
+	 type from the f16 storage type (unsigned ints).  */
+      expr = add_temp_var ("wide_cmp_result", expr);
+      tree_stl_vec wide_elements;
+      tree_stl_vec shrunk_elements;
+      unpack (expr, wide_elements);
+      for (size_t i = 0; i < wide_elements.size (); ++i)
+	{
+	  tree wide = wide_elements.at (i);
+	  shrunk_elements.push_back
+	    (convert_to_integer (short_integer_type_node, wide));
+	}
+      expr = pack (shrunk_elements);
+    }
+  build_output_assignment (*inst_base, operands[0], expr);
+
+  return base->byteCount;
+}

gcc/brig/brigfrontend/brig-comment-handler.cc

+/* brig-comment-handler.cc -- brig comment directive handling
+   Copyright (C) 2016 Free Software Foundation, Inc.
+   Contributed by Pekka Jaaskelainen <pekka.jaaskelainen@parmance.com>
+   for General Processor Tech.
+
+This file is part of GCC.
+
+GCC is free software; you can redistribute it and/or modify it under
+the terms of the GNU General Public License as published by the Free
+Software Foundation; either version 3, or (at your option) any later
+version.
+
+GCC is distributed in the hope that it will be useful, but WITHOUT ANY
+WARRANTY; without even the implied warranty of MERCHANTABILITY or
+FITNESS FOR A PARTICULAR PURPOSE.  See the GNU General Public License
+for more details.
+
+You should have received a copy of the GNU General Public License
+along with GCC; see the file COPYING3.  If not see
+<http://www.gnu.org/licenses/>.  */
+
+#include "brig-code-entry-handler.h"
+
+extern int gccbrig_verbose;
+
+size_t
+brig_directive_comment_handler::operator () (const BrigBase *base)
+{
+  const BrigDirectiveComment *brig_comment
+    = (const BrigDirectiveComment *) base;
+
+  if (gccbrig_verbose)
+    {
+      std::string cmnt = m_parent.get_string (brig_comment->name);
+      fprintf (stderr, "brig: Comment: '%s'\n", cmnt.c_str());
+    }
+  return base->byteCount;
+}

gcc/brig/brigfrontend/brig-control-handler.cc

+/* brig-control-handler.cc -- brig control directive handling
+   Copyright (C) 2016 Free Software Foundation, Inc.
+   Contributed by Pekka Jaaskelainen <pekka.jaaskelainen@parmance.com>
+   for General Processor Tech.
+
+This file is part of GCC.
+
+GCC is free software; you can redistribute it and/or modify it under
+the terms of the GNU General Public License as published by the Free
+Software Foundation; either version 3, or (at your option) any later
+version.
+
+GCC is distributed in the hope that it will be useful, but WITHOUT ANY
+WARRANTY; without even the implied warranty of MERCHANTABILITY or
+FITNESS FOR A PARTICULAR PURPOSE.  See the GNU General Public License
+for more details.
+
+You should have received a copy of the GNU General Public License
+along with GCC; see the file COPYING3.  If not see
+<http://www.gnu.org/licenses/>.  */
+
+#include "brig-code-entry-handler.h"
+#include "diagnostic.h"
+#include "print-tree.h"
+
+size_t
+brig_directive_control_handler::operator () (const BrigBase *base)
+{
+  const BrigDirectiveControl *inst = (const BrigDirectiveControl *) base;
+  const BrigData *operand_entries
+    = m_parent.get_brig_data_entry (inst->operands);
+
+  /* Parse the constant integer operands.  */
+  std::vector<tree> operands;
+  for (size_t i = 0; i < operand_entries->byteCount / 4; ++i)
+    {
+      uint32_t operand_offset
+	= ((const uint32_t *) &operand_entries->bytes)[i];
+      const BrigBase *operand_data
+	= m_parent.get_brig_operand_entry (operand_offset);
+
+      tree operand_type
+	= (inst->control == BRIG_CONTROL_REQUIREDGRIDSIZE
+	   || inst->control == BRIG_CONTROL_MAXFLATGRIDSIZE) ?
+	uint64_type_node : uint32_type_node;
+      operands.push_back
+	(build_tree_operand (*(const BrigInstBase*)inst, *operand_data,
+			     operand_type));
+    }
+
+  switch (inst->control)
+    {
+    case BRIG_CONTROL_MAXDYNAMICGROUPSIZE:
+      {
+	m_parent.m_cf->m_descriptor.max_dynamic_group_size
+	  = int_constant_value (operands.at (0));
+	break;
+      }
+    case BRIG_CONTROL_MAXFLATGRIDSIZE:
+      {
+	m_parent.m_cf->m_descriptor.max_flat_grid_size
+	  = int_constant_value (operands.at (0));
+	break;
+      }
+    case BRIG_CONTROL_MAXFLATWORKGROUPSIZE:
+      {
+	m_parent.m_cf->m_descriptor.max_flat_workgroup_size
+	  = int_constant_value (operands.at (0));
+	break;
+      }
+    case BRIG_CONTROL_REQUIREDDIM:
+      {
+	m_parent.m_cf->m_descriptor.required_dim
+	  = int_constant_value (operands.at (0));
+	break;
+      }
+    case BRIG_CONTROL_REQUIREDGRIDSIZE:
+      {
+	m_parent.m_cf->m_descriptor.required_grid_size[0]
+	  = int_constant_value (operands.at (0));
+	m_parent.m_cf->m_descriptor.required_grid_size[1]
+	  = int_constant_value (operands.at (1));
+	m_parent.m_cf->m_descriptor.required_grid_size[2]
+	  = int_constant_value (operands.at (2));
+	break;
+      }
+    case BRIG_CONTROL_REQUIREDWORKGROUPSIZE:
+      {
+	m_parent.m_cf->m_descriptor.required_workgroup_size[0]
+	  = int_constant_value (operands.at (0));
+	m_parent.m_cf->m_descriptor.required_workgroup_size[1]
+	  = int_constant_value (operands.at (1));
+	m_parent.m_cf->m_descriptor.required_workgroup_size[2]
+	  = int_constant_value (operands.at (2));
+	break;
+      }
+    case BRIG_CONTROL_REQUIRENOPARTIALWORKGROUPS:
+      /* Performance hint only, ignored for now.  */
+      break;
+    case BRIG_CONTROL_ENABLEBREAKEXCEPTIONS:
+    case BRIG_CONTROL_ENABLEDETECTEXCEPTIONS:
+      /* Unimplemented.  */
+      break;
+    default:
+      sorry ("Unsupported control directive %x.\n", inst->control);
+    }
+  return base->byteCount;
+}

gcc/brig/brigfrontend/brig-copy-move-inst-handler.cc

+/* brig-copy-move-inst-handler.cc -- brig copy/move instruction handling
+   Copyright (C) 2016 Free Software Foundation, Inc.
+   Contributed by Pekka Jaaskelainen <pekka.jaaskelainen@parmance.com>
+   for General Processor Tech.
+
+   This file is part of GCC.
+
+   GCC is free software; you can redistribute it and/or modify it under
+   the terms of the GNU General Public License as published by the Free
+   Software Foundation; either version 3, or (at your option) any later
+   version.
+
+   GCC is distributed in the hope that it will be useful, but WITHOUT ANY
+   WARRANTY; without even the implied warranty of MERCHANTABILITY or
+   FITNESS FOR A PARTICULAR PURPOSE.  See the GNU General Public License
+   for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with GCC; see the file COPYING3.  If not see
+   <http://www.gnu.org/licenses/>.  */
+
+#include "brig-code-entry-handler.h"
+#include "tree-pretty-print.h"
+#include "print-tree.h"
+#include "errors.h"
+#include "brig-util.h"
+
+size_t
+brig_copy_move_inst_handler::handle_lda (const BrigInstBase *brig_inst)
+{
+  tree dest_type = gccbrig_tree_type_for_hsa_type (brig_inst->type);
+
+  tree input = build_tree_operand_from_brig (brig_inst, NULL, 1);
+  tree output = build_tree_operand_from_brig (brig_inst, dest_type, 0);
+
+  build_output_assignment (*brig_inst, output, input);
+  return brig_inst->base.byteCount;
+}
+
+size_t
+brig_copy_move_inst_handler::operator () (const BrigBase *base)
+{
+  const BrigInstBase *brig_inst
+    = (const BrigInstBase *) &((const BrigInstBasic *) base)->base;
+
+  if (brig_inst->opcode == BRIG_OPCODE_LDA)
+    return handle_lda (brig_inst);
+
+  const BrigInstSourceType *inst_src_type = (const BrigInstSourceType *) base;
+
+  tree source_type = gccbrig_tree_type_for_hsa_type (inst_src_type->sourceType);
+  tree dest_type = gccbrig_tree_type_for_hsa_type (brig_inst->type);
+
+  tree input = build_tree_operand_from_brig (brig_inst, source_type, 1);
+  tree output = build_tree_operand_from_brig (brig_inst, dest_type, 0);
+  if (brig_inst->opcode == BRIG_OPCODE_COMBINE)
+    {
+      /* For combine, a simple reinterpret cast from the array constructor
+	 works.  */
+
+      tree casted = build_reinterpret_cast (dest_type, input);
+      tree assign = build2 (MODIFY_EXPR, TREE_TYPE (output), output, casted);
+      m_parent.m_cf->append_statement (assign);
+    }
+  else if (brig_inst->opcode == BRIG_OPCODE_EXPAND)
+    build_output_assignment (*brig_inst, output, input);
+  else
+    {
+      brig_basic_inst_handler basic (m_parent);
+      return basic (base);
+    }
+  return base->byteCount;
+}

gcc/brig/brigfrontend/brig-cvt-inst-handler.cc

+/* brig-cvt-inst-handler.cc -- brig cvt (convert) instruction handling
+   Copyright (C) 2016 Free Software Foundation, Inc.
+   Contributed by Pekka Jaaskelainen <pekka.jaaskelainen@parmance.com>
+   for General Processor Tech.
+
+   This file is part of GCC.
+
+   GCC is free software; you can redistribute it and/or modify it under
+   the terms of the GNU General Public License as published by the Free
+   Software Foundation; either version 3, or (at your option) any later
+   version.
+
+   GCC is distributed in the hope that it will be useful, but WITHOUT ANY
+   WARRANTY; without even the implied warranty of MERCHANTABILITY or
+   FITNESS FOR A PARTICULAR PURPOSE.  See the GNU General Public License
+   for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with GCC; see the file COPYING3.  If not see
+   <http://www.gnu.org/licenses/>.  */
+
+#include <sstream>
+
+#include "brig-code-entry-handler.h"
+
+#include "gimple-expr.h"
+#include "errors.h"
+#include "convert.h"
+#include "tree-pretty-print.h"
+#include "print-tree.h"
+#include "diagnostic-core.h"
+#include "brig-util.h"
+
+const BrigAluModifier8_t *
+brig_cvt_inst_handler::modifier (const BrigBase *base) const
+{
+  const BrigInstCvt *inst = (const BrigInstCvt *) base;
+  return &inst->modifier;
+}
+
+const BrigRound8_t *
+brig_cvt_inst_handler::round (const BrigBase *base) const
+{
+  const BrigInstCvt *inst = (const BrigInstCvt *) base;
+  return &inst->round;
+}
+
+size_t
+brig_cvt_inst_handler::generate (const BrigBase *base)
+{
+  /* In cvt instructions there can be at least four data types involved:
+
+     - the input register type
+     - the output register type
+     - the conversion source type
+     - the conversion destination type
+  */
+
+  const BrigInstBase *brig_inst
+    = (const BrigInstBase *) &((const BrigInstBasic *) base)->base;
+  const BrigInstCvt *cvt_inst = (const BrigInstCvt *) base;
+
+  const BrigAluModifier8_t *inst_modifier = modifier (base);
+  const bool FTZ = inst_modifier != NULL && (*inst_modifier) & BRIG_ALU_FTZ;
+
+  /* The conversion source type.  */
+  tree src_type = get_tree_expr_type_for_hsa_type (cvt_inst->sourceType);
+
+  bool src_is_fp16 = cvt_inst->sourceType == BRIG_TYPE_F16;
+
+  /* The conversion destination type.  */
+  tree dest_type = gccbrig_tree_type_for_hsa_type (brig_inst->type);
+
+  bool dest_is_fp16 = brig_inst->type == BRIG_TYPE_F16;
+
+  if (!dest_type || !src_type)
+    {
+      gcc_unreachable ();
+      return base->byteCount;
+    }
+
+  tree_stl_vec operands = build_operands (*brig_inst);
+  tree &input = operands.at (1);
+  tree &output = operands.at (0);
+
+  size_t conv_src_size = int_size_in_bytes (src_type);
+  size_t conv_dst_size = int_size_in_bytes (dest_type);
+  size_t src_reg_size = int_size_in_bytes (TREE_TYPE (input));
+
+  /* The input register can be of different type&size than the
+     conversion input size.  First cast the input to the conversion
+     input type.  These casts are always bitcasts which can be
+     expressed as casts between different unsigned integers.  */
+  if (src_reg_size != conv_src_size)
+    {
+      tree unsigned_int_type = NULL_TREE;
+      if (INTEGRAL_TYPE_P (src_type))
+	unsigned_int_type = unsigned_type_for (src_type);
+      else /* Find a matching size int type for the REAL type.  */
+	{
+	  if (conv_src_size == 2)
+	    unsigned_int_type = gccbrig_tree_type_for_hsa_type (BRIG_TYPE_U16);
+	  else if (conv_src_size == 4)
+	    unsigned_int_type = gccbrig_tree_type_for_hsa_type (BRIG_TYPE_U32);
+	  else if (conv_src_size == 8)
+	    unsigned_int_type = gccbrig_tree_type_for_hsa_type (BRIG_TYPE_U64);
+	  else
+	    gcc_unreachable ();
+	}
+      input = convert_to_integer (unsigned_int_type, input);
+    }
+
+  if (src_is_fp16)
+    input = build_h2f_conversion (input);
+
+  /* Flush the float operand to zero if indicated with 'ftz'.  */
+  if (FTZ && SCALAR_FLOAT_TYPE_P (src_type))
+    {
+      tree casted_input = build_reinterpret_cast (src_type, input);
+      input = flush_to_zero (src_is_fp16) (*this, casted_input);
+    }
+
+  tree conversion_result = NULL_TREE;
+  if (brig_inst->type == BRIG_TYPE_B1)
+    {
+      /* When the destination is b1, cvt does a 'ztest' operation which is
+	 defined as a != 0 for integers and similarly (!= 0.0f) for floats.  */
+      if (INTEGRAL_TYPE_P (src_type))
+	{
+	  /* Generate an integer not equal operation.  */
+	  conversion_result = build2 (NE_EXPR, TREE_TYPE (input), input,
+				      build_int_cst (TREE_TYPE (input), 0));
+	}
+      else
+	{
+	  /* For REAL source types, ztest returns 1 if the value is not +- 0.0f.
+	     We can perform this check with an integer comparison after
+	     masking away the sign bit from a correct position.  This is safer
+	     than using absf because of exceptions in case of a NaN
+	     input (NaN exceptions are not generated with cvt).  */
+	  tree unsigned_int_type = NULL_TREE;
+	  /* Bit battern with all but the upper bit 1.  */
+	  tree and_mask = NULL_TREE;
+	  if (conv_src_size == 2)
+	    {
+	      unsigned_int_type = gccbrig_tree_type_for_hsa_type (BRIG_TYPE_U16);
+	      and_mask = build_int_cst (unsigned_int_type, 0x7FFF);
+	    }
+	  else if (conv_src_size == 4)
+	    {
+	      unsigned_int_type = gccbrig_tree_type_for_hsa_type (BRIG_TYPE_U32);
+	      and_mask = build_int_cst (unsigned_int_type, 0x7FFFFFFF);
+	    }
+	  else if (conv_src_size == 8)
+	    {
+	      unsigned_int_type = gccbrig_tree_type_for_hsa_type (BRIG_TYPE_U64);
+	      and_mask = build_int_cst (unsigned_int_type, 0x7FFFFFFFFFFFFFFF);
+	    }
+	  else
+	    gcc_unreachable ();
+	  tree casted_input = build_reinterpret_cast (unsigned_int_type, input);
+	  tree masked_input
+	    = build2 (BIT_AND_EXPR, unsigned_int_type, casted_input, and_mask);
+	  conversion_result
+	    = build2 (NE_EXPR, TREE_TYPE (masked_input), masked_input,
+		      build_int_cst (unsigned_int_type, 0));
+	}
+      /* The result from the comparison is a boolean, convert it to such.  */
+      conversion_result
+	= convert_to_integer (gccbrig_tree_type_for_hsa_type (BRIG_TYPE_B1),
+			      conversion_result);
+    }
+  else if (dest_is_fp16)
+    {
+      tree casted_input = build_reinterpret_cast (src_type, input);
+      conversion_result
+	= convert_to_real (brig_to_generic::s_fp32_type, casted_input);
+      if (FTZ)
+	conversion_result = flush_to_zero (true) (*this, conversion_result);
+      conversion_result = build_f2h_conversion (conversion_result);
+    }
+  else if (SCALAR_FLOAT_TYPE_P (dest_type))
+    {
+      tree casted_input = build_reinterpret_cast (src_type, input);
+      conversion_result = convert_to_real (dest_type, casted_input);
+    }
+  else if (INTEGRAL_TYPE_P (dest_type) && INTEGRAL_TYPE_P (src_type))
+    {
+      conversion_result = extend_int (input, dest_type, src_type);
+    }
+  else if (INTEGRAL_TYPE_P (dest_type) && SCALAR_FLOAT_TYPE_P (src_type))
+    {
+
+      if (cvt_inst->round == BRIG_ROUND_INTEGER_ZERO_SAT)
+	{
+
+	  /* Use builtins for the saturating conversions.  */
+#undef DEF_HSAIL_SAT_BUILTIN
+#undef DEF_HSAIL_BUILTIN
+#undef DEF_HSAIL_ATOMIC_BUILTIN
+#undef DEF_HSAIL_INTR_BUILTIN
+#undef DEF_HSAIL_CVT_ZEROI_SAT_BUILTIN
+
+	  tree builtin = NULL_TREE;
+	  BrigType16_t src_arith_type
+	    = src_is_fp16
+	    ? (BrigType16_t) BRIG_TYPE_F32 : cvt_inst->sourceType;
+#define DEF_HSAIL_CVT_ZEROI_SAT_BUILTIN(ENUM, HSAIL_DST_TYPE, HSAIL_SRC_TYPE, \
+					NAME, TYPE, ATTRS)		\
+	  if (brig_inst->type == HSAIL_DST_TYPE				\
+	      && src_arith_type == HSAIL_SRC_TYPE)		\
+	    builtin = builtin_decl_explicit (ENUM);			\
+	  else
+#include "brig-builtins.def"
+	    gcc_unreachable ();
+
+	  tree casted_input = build_reinterpret_cast (src_type, input);
+	  conversion_result
+	    = call_builtin (builtin, 1, dest_type, src_type, casted_input);
+	}
+      else
+	{
+	  tree casted_input = build_reinterpret_cast (src_type, input);
+
+	  /* Perform the int to float conversion.  */
+	  conversion_result = convert_to_integer (dest_type, casted_input);
+	}
+      /* The converted result is finally extended to the target register
+	 width, using the same sign as the destination.  */
+      conversion_result
+	= convert_to_integer (TREE_TYPE (output), conversion_result);
+    }
+  else
+    {
+      /* Just use CONVERT_EXPR and hope for the best.  */
+      tree casted_input = build_reinterpret_cast (dest_type, input);
+      conversion_result = build1 (CONVERT_EXPR, dest_type, casted_input);
+    }
+
+  size_t dst_reg_size = int_size_in_bytes (TREE_TYPE (output));
+
+  tree assign = NULL_TREE;
+  /* The output register can be of different type&size than the
+     conversion output size.  Cast it to the register variable type.  */
+  if (dst_reg_size > conv_dst_size)
+    {
+      tree casted_output
+	= build1 (CONVERT_EXPR, TREE_TYPE (output), conversion_result);
+      assign = build2 (MODIFY_EXPR, TREE_TYPE (output), output, casted_output);
+    }
+  else
+    {
+      tree casted_output
+	= build_reinterpret_cast (TREE_TYPE (output), conversion_result);
+      assign = build2 (MODIFY_EXPR, TREE_TYPE (output), output, casted_output);
+    }
+  m_parent.m_cf->append_statement (assign);
+
+  return base->byteCount;
+}

gcc/brig/brigfrontend/brig-fbarrier-handler.cc

+/* brig-fbarrier-handler.cc -- brig fbarrier directive handling
+   Copyright (C) 2016 Free Software Foundation, Inc.
+   Contributed by Pekka Jaaskelainen <pekka.jaaskelainen@parmance.com>
+   for General Processor Tech.
+
+This file is part of GCC.
+
+GCC is free software; you can redistribute it and/or modify it under
+the terms of the GNU General Public License as published by the Free
+Software Foundation; either version 3, or (at your option) any later
+version.
+
+GCC is distributed in the hope that it will be useful, but WITHOUT ANY
+WARRANTY; without even the implied warranty of MERCHANTABILITY or
+FITNESS FOR A PARTICULAR PURPOSE.  See the GNU General Public License
+for more details.
+
+You should have received a copy of the GNU General Public License
+along with GCC; see the file COPYING3.  If not see
+<http://www.gnu.org/licenses/>.  */
+
+#include "brig-code-entry-handler.h"
+
+#include "stringpool.h"
+#include "errors.h"
+
+/* Allocate this many bytes from the group segment for each fbarrier.  */
+#define FBARRIER_STRUCT_SIZE 32
+
+size_t
+brig_directive_fbarrier_handler::operator () (const BrigBase *base)
+{
+  /* Model fbarriers as group segment variables with fixed size
+     large enough to store whatever data the actual target needs
+     to store to maintain the barrier info.  The handle is the
+     offset to the beginning of the object.  */
+
+  const BrigDirectiveFbarrier* fbar = (const BrigDirectiveFbarrier*)base;
+  if (m_parent.m_cf != NULL)
+    m_parent.m_cf->m_function_scope_vars.insert (base);
+  std::string var_name = m_parent.get_mangled_name (fbar);
+  m_parent.append_group_variable (var_name, FBARRIER_STRUCT_SIZE, 1);
+  return base->byteCount;
+}

gcc/brig/brigfrontend/brig-function.h

+/* brig-function.h -- declaration of brig_function class.
+   Copyright (C) 2016 Free Software Foundation, Inc.
+   Contributed by Pekka Jaaskelainen <pekka.jaaskelainen@parmance.com>
+   for General Processor Tech.
+
+   This file is part of GCC.
+
+   GCC is free software; you can redistribute it and/or modify it under
+   the terms of the GNU General Public License as published by the Free
+   Software Foundation; either version 3, or (at your option) any later
+   version.
+
+   GCC is distributed in the hope that it will be useful, but WITHOUT ANY
+   WARRANTY; without even the implied warranty of MERCHANTABILITY or
+   FITNESS FOR A PARTICULAR PURPOSE.  See the GNU General Public License
+   for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with GCC; see the file COPYING3.  If not see
+   <http://www.gnu.org/licenses/>.  */
+
+#ifndef BRIG_FUNCTION_H
+#define BRIG_FUNCTION_H
+
+#include "config.h"
+#include "system.h"
+#include "ansidecl.h"
+#include "coretypes.h"
+#include "opts.h"
+#include "tree.h"
+#include "tree-iterator.h"
+#include "hsa-brig-format.h"
+
+class brig_to_generic;
+
+#include <map>
+#include <string>
+#include <vector>
+#include <set>
+
+#include "phsa.h"
+
+typedef std::map<std::string, tree> label_index;
+typedef std::map<const BrigDirectiveVariable *, tree> variable_index;
+typedef std::vector<tree> tree_stl_vec;
+
+/* There are 128 c regs and 2048 s/d/q regs each in the HSAIL.  */
+#define BRIG_2_TREE_HSAIL_C_REG_COUNT (128)
+#define BRIG_2_TREE_HSAIL_S_REG_COUNT (2048)
+#define BRIG_2_TREE_HSAIL_D_REG_COUNT (2048)
+#define BRIG_2_TREE_HSAIL_Q_REG_COUNT (2048)
+#define BRIG_2_TREE_HSAIL_TOTAL_REG_COUNT				       \
+  (BRIG_2_TREE_HSAIL_C_REG_COUNT + BRIG_2_TREE_HSAIL_S_REG_COUNT	       \
+   + BRIG_2_TREE_HSAIL_D_REG_COUNT + BRIG_2_TREE_HSAIL_Q_REG_COUNT)
+
+/* Holds data for the currently built GENERIC function.  */
+
+class brig_function
+{
+public:
+  typedef std::map<const BrigDirectiveVariable *, size_t> var_offset_table;
+
+private:
+  struct reg_decl_index_entry
+  {
+    tree m_var_decl;
+  };
+
+public:
+  brig_function (const BrigDirectiveExecutable *exec, brig_to_generic *parent);
+  ~brig_function ();
+
+  tree arg_variable (const BrigDirectiveVariable *var) const;
+  void add_arg_variable (const BrigDirectiveVariable *brigVar, tree treeDecl);
+
+  void append_kernel_arg (const BrigDirectiveVariable *var, size_t size,
+			  size_t alignment);
+
+  size_t kernel_arg_offset (const BrigDirectiveVariable *var) const;
+
+  void add_id_variables ();
+
+  tree label (const std::string &name);
+
+  tree add_local_variable (std::string name, tree type);
+
+  tree get_m_var_declfor_reg (const BrigOperandRegister *reg);
+
+  bool convert_to_wg_function ();
+
+  void add_wi_loop (int dim, tree_stmt_iterator *header_entry,
+		    tree_stmt_iterator *branch_after);
+
+  tree emit_metadata (tree stmt_list);
+  tree emit_launcher_and_metadata ();
+
+  tree append_statement (tree stmt);
+
+  void create_alloca_frame ();
+
+  void finish ();
+  void finish_kernel ();
+
+  void append_return_stmt ();
+
+  bool has_function_scope_var (const BrigBase* var) const;
+
+  void analyze_calls ();
+
+  const BrigDirectiveExecutable *m_brig_def;
+
+  bool m_is_kernel;
+  bool m_is_finished;
+  std::string m_name;
+  tree m_current_bind_expr;
+  tree m_func_decl;
+  tree m_entry_label_stmt;
+  tree m_exit_label;
+
+  /* The __context function argument.  */
+  tree m_context_arg;
+  /* The __group_base_ptr argument in the current function.
+     Points to the start of the group segment for the kernel
+     instance.  */
+  tree m_group_base_arg;
+  /* The __private_base_ptr argument in the current function.
+     Points to the start of the private segment.  */
+  tree m_private_base_arg;
+
+  /* The return value variable for the current function.  */
+  tree m_ret_value;
+
+  /* The offsets of the kernel arguments in the __arg blob
+     pointing to the kernel argument space.  */
+  size_t m_next_kernarg_offset;
+
+  /* The largest kernel argument variable alignment.  */
+  size_t m_kernarg_max_align;
+
+  var_offset_table m_kernarg_offsets;
+
+  /* Argument variables in the currently handled binding expression
+     (argument segment).  */
+  variable_index m_arg_variables;
+
+  /* The brig variable for the function return value.  */
+  const BrigDirectiveVariable *m_ret_value_brig_var;
+
+  /* The function local temporary variable for the return value.  */
+  tree m_ret_temp;
+
+  /* Labels in the current function are collected here so we can refer
+     to them from jumps before they have been placed to the function.  */
+  label_index m_label_index;
+
+  /* If the kernel contains at least one barrier, this is set to true.  */
+  bool m_has_barriers;
+
+  /* True if the function has at least one alloca instruction.  */
+  bool m_has_allocas;
+
+  /* If the kernel containts at least one function call that _may_
+     contain a barrier call, this is set to true.  */
+  bool m_has_function_calls_with_barriers;
+
+  /* Set to true after this function has been analyzed for barrier and
+     dispatch packet instruction usage in the final call graph analysis.  */
+  bool m_calls_analyzed;
+
+  /* True in case the function was successfully converted to a WG function.  */
+  bool m_is_wg_function;
+
+  /* Work-item ID related variables are cached in the entry of the kernel
+     function in order to use them directly in address computations, leading
+     to more efficient optimizations.  The references to the local variables
+     are stored here.  */
+  tree m_local_id_vars[3];
+  tree m_cur_wg_size_vars[3];
+  tree m_wg_id_vars[3];
+  tree m_wg_size_vars[3];
+  tree m_grid_size_vars[3];
+
+  /* Set to true in case the kernel contains at least one dispatch packet
+     (work-item ID-related) builtin call that could not be expanded to
+     tree nodes.  */
+  bool m_has_unexpanded_dp_builtins;
+
+  /* Points to the instruction after which the real kernel code starts.
+     Usually points to the last WI ID variable initialization statement.  */
+  tree_stmt_iterator m_kernel_entry;
+
+  /* True if we are currently generating the contents of an arg block.  */
+  bool m_generating_arg_block;
+
+  /* A collection of function scope variables seen so far for resolving
+     variable references vs. module scope declarations.  */
+  std::set<const BrigBase*> m_function_scope_vars;
+
+  /* The functions called by this function.  */
+  std::vector<tree> m_called_functions;
+
+  brig_to_generic *m_parent;
+  /* The metadata of the function that should be stored with the binary and
+     passed to the HSA runtime:  */
+  phsa_descriptor m_descriptor;
+
+private:
+  /* Bookkeeping for the different HSA registers and their tree declarations
+     for the currently generated function.  */
+  reg_decl_index_entry *m_regs[BRIG_2_TREE_HSAIL_TOTAL_REG_COUNT];
+};
+
+#endif

gcc/brig/brigfrontend/brig-inst-mod-handler.cc

+/* brig-inst-mod-handler.cc -- brig rounding moded instruction handling
+   Copyright (C) 2016 Free Software Foundation, Inc.
+   Contributed by Pekka Jaaskelainen <pekka.jaaskelainen@parmance.com>
+   for General Processor Tech.
+
+This file is part of GCC.
+
+GCC is free software; you can redistribute it and/or modify it under
+the terms of the GNU General Public License as published by the Free
+Software Foundation; either version 3, or (at your option) any later
+version.
+
+GCC is distributed in the hope that it will be useful, but WITHOUT ANY
+WARRANTY; without even the implied warranty of MERCHANTABILITY or
+FITNESS FOR A PARTICULAR PURPOSE.  See the GNU General Public License
+for more details.
+
+You should have received a copy of the GNU General Public License
+along with GCC; see the file COPYING3.  If not see
+<http://www.gnu.org/licenses/>.  */
+
+#include "brig-code-entry-handler.h"
+
+#include "gimple-expr.h"
+#include "errors.h"
+
+size_t
+brig_inst_mod_handler::generate (const BrigBase *base)
+{
+  brig_basic_inst_handler basic_handler (m_parent);
+  return basic_handler (base);
+}
+
+const BrigAluModifier8_t *
+brig_inst_mod_handler::modifier (const BrigBase *base) const
+{
+  const BrigInstMod *inst = (const BrigInstMod *) base;
+  return &inst->modifier;
+}
+
+const BrigRound8_t *
+brig_inst_mod_handler::round (const BrigBase *base) const
+{
+  const BrigInstMod *inst = (const BrigInstMod *) base;
+  return &inst->round;
+}
+
+/* This used to inject fesetround () calls to control the rounding mode of the
+   actual executed floating point operation.  It turned out that supporting
+   conversions using fesetround calls won't work in gcc due to it not being
+   able to restrict code motions across calls at the moment.  This
+   functionality is therefore disabled for now until a better solution is
+   found or if fesetround () is fixed in gcc.  */
+size_t
+brig_inst_mod_handler::operator () (const BrigBase *base)
+{
+  return generate (base);
+}

gcc/brig/brigfrontend/brig-label-handler.cc

+/* brig-label-handler.cc -- brig label directive handling
+   Copyright (C) 2016 Free Software Foundation, Inc.
+   Contributed by Pekka Jaaskelainen <pekka.jaaskelainen@parmance.com>
+   for General Processor Tech.
+
+This file is part of GCC.
+
+GCC is free software; you can redistribute it and/or modify it under
+the terms of the GNU General Public License as published by the Free
+Software Foundation; either version 3, or (at your option) any later
+version.
+
+GCC is distributed in the hope that it will be useful, but WITHOUT ANY
+WARRANTY; without even the implied warranty of MERCHANTABILITY or
+FITNESS FOR A PARTICULAR PURPOSE.  See the GNU General Public License
+for more details.
+
+You should have received a copy of the GNU General Public License
+along with GCC; see the file COPYING3.  If not see
+<http://www.gnu.org/licenses/>.  */
+
+#include "brig-code-entry-handler.h"
+
+size_t
+brig_directive_label_handler::operator () (const BrigBase *base)
+{
+  const BrigDirectiveLabel *brig_label = (const BrigDirectiveLabel *) base;
+
+  const BrigData *label_name = m_parent.get_brig_data_entry (brig_label->name);
+
+  std::string label_str ((const char *) (label_name->bytes),
+			 label_name->byteCount);
+
+  tree stmt = build_stmt (LABEL_EXPR, m_parent.m_cf->label (label_str));
+  m_parent.m_cf->append_statement (stmt);
+  return base->byteCount;
+}

gcc/brig/brigfrontend/brig-lane-inst-handler.cc

+/* brig-lane-inst-handler.cc -- brig lane instruction handling
+   Copyright (C) 2016 Free Software Foundation, Inc.
+   Contributed by Pekka Jaaskelainen <pekka.jaaskelainen@parmance.com>
+   for General Processor Tech.
+
+This file is part of GCC.
+
+GCC is free software; you can redistribute it and/or modify it under
+the terms of the GNU General Public License as published by the Free
+Software Foundation; either version 3, or (at your option) any later
+version.
+
+GCC is distributed in the hope that it will be useful, but WITHOUT ANY
+WARRANTY; without even the implied warranty of MERCHANTABILITY or
+FITNESS FOR A PARTICULAR PURPOSE.  See the GNU General Public License
+for more details.
+
+You should have received a copy of the GNU General Public License
+along with GCC; see the file COPYING3.  If not see
+<http://www.gnu.org/licenses/>.  */
+
+#include "brig-code-entry-handler.h"
+#include "errors.h"
+#include "diagnostic-core.h"
+#include "brig-util.h"
+
+brig_lane_inst_handler::brig_lane_inst_handler (brig_to_generic &parent)
+  : brig_code_entry_handler (parent)
+{
+}
+
+size_t
+brig_lane_inst_handler::operator () (const BrigBase *base)
+{
+  const BrigInstLane &inst = *(const BrigInstLane *) base;
+  tree_stl_vec operands = build_operands (inst.base);
+
+  tree expr = NULL_TREE;
+  if (inst.base.opcode == BRIG_OPCODE_ACTIVELANECOUNT)
+    {
+      /* Because we are fixed to single WI per wave, it's enough to
+	 just check the src value of the single work item itself.  */
+      expr = build2 (NE_EXPR, uint32_type_node,
+		     build_zero_cst (uint32_type_node), operands[1]);
+    }
+  else if (inst.base.opcode == BRIG_OPCODE_ACTIVELANEID)
+    {
+      expr = build_zero_cst (uint32_type_node);
+    }
+  else if (inst.base.opcode == BRIG_OPCODE_ACTIVELANEMASK)
+    {
+      tree u64_type = gccbrig_tree_type_for_hsa_type (BRIG_TYPE_U64);
+      tree zero_cst = build_zero_cst (u64_type);
+      expr = build2 (NE_EXPR, u64_type, zero_cst, operands[1]);
+
+      tree_stl_vec elements;
+      elements.push_back (expr);
+      elements.push_back (zero_cst);
+      elements.push_back (zero_cst);
+      elements.push_back (zero_cst);
+
+      expr = pack (elements);
+    }
+  else if (inst.base.opcode == BRIG_OPCODE_ACTIVELANEPERMUTE)
+    {
+      tree src = operands[1];
+      tree identity = operands[3];
+      tree use_identity = operands[4];
+
+      /* When WAVESIZE is 1, we either select the src of the work-item
+	 itself or 'identity' in case use_identity is 1.  */
+
+      tree cmp = build2 (EQ_EXPR, uint32_type_node,
+			 build_int_cstu (uint32_type_node, 1), use_identity);
+
+      expr = build3 (COND_EXPR, TREE_TYPE (src), cmp, identity, src);
+    }
+  else
+    gcc_unreachable ();
+
+  build_output_assignment (inst.base, operands[0], expr);
+
+  return base->byteCount;
+}

gcc/brig/brigfrontend/brig-machine.c

+/* brig-machine.c -- gccbrig machine queries
+   Copyright (C) 2016 Free Software Foundation, Inc.
+   Contributed by Pekka Jaaskelainen <pekka.jaaskelainen@parmance.com>
+   for General Processor Tech.
+
+This file is part of GCC.
+
+GCC is free software; you can redistribute it and/or modify it under
+the terms of the GNU General Public License as published by the Free
+Software Foundation; either version 3, or (at your option) any later
+version.
+
+GCC is distributed in the hope that it will be useful, but WITHOUT ANY
+WARRANTY; without even the implied warranty of MERCHANTABILITY or
+FITNESS FOR A PARTICULAR PURPOSE.  See the GNU General Public License
+for more details.
+
+You should have received a copy of the GNU General Public License
+along with GCC; see the file COPYING3.  If not see
+<http://www.gnu.org/licenses/>.  */
+
+#include "config.h"
+#include "system.h"
+#include "coretypes.h"
+#include "brig-machine.h"
+
+/* Return the numerical address space id for the segment in the current
+   target.  Currently a dummy function that always returns 0, serves as
+   a placeholder for multi-AS machines.  */
+
+unsigned
+gccbrig_get_target_addr_space_id (BrigSegment8_t)
+{
+  return 0;
+}
+
+/* Return the WAVESIZE for the current target.  For now a dummy placeholder
+   returning always 1.  */
+
+unsigned
+gccbrig_get_target_wavesize ()
+{
+  return 1;
+}

gcc/brig/brigfrontend/brig-machine.h

+/* brig-machine.h -- gccbrig machine queries
+   Copyright (C) 2016 Free Software Foundation, Inc.
+
+This file is part of GCC.
+
+GCC is free software; you can redistribute it and/or modify it under
+the terms of the GNU General Public License as published by the Free
+Software Foundation; either version 3, or (at your option) any later
+version.
+
+GCC is distributed in the hope that it will be useful, but WITHOUT ANY
+WARRANTY; without even the implied warranty of MERCHANTABILITY or
+FITNESS FOR A PARTICULAR PURPOSE.  See the GNU General Public License
+for more details.
+
+You should have received a copy of the GNU General Public License
+along with GCC; see the file COPYING3.  If not see
+<http://www.gnu.org/licenses/>.  */
+
+#ifndef GCC_BRIG_MACHINE_H
+#define GCC_BRIG_MACHINE_H
+
+#include "hsa-brig-format.h"
+
+/* These functions should be eventually converted to machine info queries and
+   redefined at backends.  At that point make these functions delegate to
+   those.  */
+
+unsigned gccbrig_get_target_addr_space_id (BrigSegment8_t segment);
+
+unsigned gccbrig_get_target_wavesize ();
+
+#endif

gcc/brig/brigfrontend/brig-mem-inst-handler.cc

+/* brig-mem-inst-handler.cc -- brig memory inst handler
+   Copyright (C) 2016 Free Software Foundation, Inc.
+   Contributed by Pekka Jaaskelainen <pekka.jaaskelainen@parmance.com>
+   for General Processor Tech.
+
+   This file is part of GCC.
+
+   GCC is free software; you can redistribute it and/or modify it under
+   the terms of the GNU General Public License as published by the Free
+   Software Foundation; either version 3, or (at your option) any later
+   version.
+
+   GCC is distributed in the hope that it will be useful, but WITHOUT ANY
+   WARRANTY; without even the implied warranty of MERCHANTABILITY or
+   FITNESS FOR A PARTICULAR PURPOSE.  See the GNU General Public License
+   for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with GCC; see the file COPYING3.  If not see
+   <http://www.gnu.org/licenses/>.  */
+
+#include "brig-code-entry-handler.h"
+
+#include "errors.h"
+#include "brig-util.h"
+#include "gimple-expr.h"
+#include "print-tree.h"
+#include "tree-pretty-print.h"
+#include "convert.h"
+#include "diagnostic-core.h"
+
+tree
+brig_mem_inst_handler::build_mem_access (const BrigInstBase *brig_inst,
+					 tree addr, tree data)
+{
+  bool is_load = brig_inst->opcode == BRIG_OPCODE_LD;
+  bool is_store = brig_inst->opcode == BRIG_OPCODE_ST;
+
+  if (!is_load && !is_store)
+    gcc_unreachable ();
+
+  tree instr_type = gccbrig_tree_type_for_hsa_type (brig_inst->type);
+
+  if (VECTOR_TYPE_P (TREE_TYPE (data)))
+    instr_type = TREE_TYPE (data);
+
+  tree ptype = build_pointer_type (instr_type);
+
+  /* The HSAIL mem instructions are unaligned by default.
+     TODO: exploit the align modifier, it should lead to faster code.
+  */
+  tree unaligned_type = build_aligned_type (instr_type, 8);
+
+  /* Create a mem ref from the previous result, without offset.  */
+  tree mem_ref
+    = build2 (MEM_REF, unaligned_type, addr, build_int_cst (ptype, 0));
+
+  if (is_load)
+    {
+      /* Add a temporary variable so there won't be multiple
+	 reads in case of vector unpack.  */
+      mem_ref = add_temp_var ("mem_read", mem_ref);
+      return build_output_assignment (*brig_inst, data, mem_ref);
+    }
+  else
+    {
+      tree stmt = build2 (MODIFY_EXPR, TREE_TYPE (mem_ref), mem_ref, data);
+      return m_parent.m_cf->append_statement (stmt);
+    }
+  return mem_ref;
+}
+
+size_t
+brig_mem_inst_handler::operator () (const BrigBase *base)
+{
+  const BrigInstBase *brig_inst
+    = (const BrigInstBase *) &((const BrigInstBasic *) base)->base;
+
+  if (brig_inst->opcode == BRIG_OPCODE_ALLOCA)
+    {
+      tree_stl_vec operands = build_operands (*brig_inst);
+      size_t alignment = 1;
+      const BrigInstMem *mem_inst = (const BrigInstMem *) brig_inst;
+      if (mem_inst->align != BRIG_ALIGNMENT_NONE)
+	{
+	  alignment = 1 << (mem_inst->align - 1);
+	}
+
+      tree align_opr = build_int_cstu (size_type_node, alignment);
+      tree_stl_vec inputs;
+      inputs.push_back (operands[1]);
+      inputs.push_back (align_opr);
+      tree builtin_call
+	= expand_or_call_builtin (BRIG_OPCODE_ALLOCA, BRIG_TYPE_U32,
+				  uint32_type_node, inputs);
+      build_output_assignment (*brig_inst, operands[0], builtin_call);
+      m_parent.m_cf->m_has_allocas = true;
+      return base->byteCount;
+    }
+
+  tree instr_type = gccbrig_tree_type_for_hsa_type (brig_inst->type);
+
+  const BrigData *operand_entries
+    = m_parent.get_brig_data_entry (brig_inst->operands);
+
+  uint32_t data_operand_offset;
+  memcpy (&data_operand_offset, &operand_entries->bytes, 4);
+
+  const BrigBase *operand
+    = m_parent.get_brig_operand_entry (data_operand_offset);
+
+  const BrigData *operandData = NULL;
+
+  bool is_store = brig_inst->opcode == BRIG_OPCODE_ST;
+
+  bool is_three_element_vector_access
+    = operand->kind == BRIG_KIND_OPERAND_OPERAND_LIST
+      && (operandData = m_parent.get_brig_data_entry
+	  (((const BrigOperandOperandList *) operand)->elements))
+      && operandData->byteCount / 4 == 3;
+
+  if (is_three_element_vector_access)
+    {
+      /* We need to scalarize the 3-element vector accesses here
+	 because gcc assumes the GENERIC vector datatypes are of two exponent
+	 size internally.  */
+      size_t bytes = operandData->byteCount;
+      const BrigOperandOffset32_t *operand_ptr
+	= (const BrigOperandOffset32_t *) operandData->bytes;
+
+      uint32_t addr_operand_offset;
+      memcpy (&addr_operand_offset, &operand_entries->bytes + 4, 4);
+
+      const BrigOperandAddress *addr_operand
+	= (const BrigOperandAddress *) m_parent.get_brig_operand_entry
+	(addr_operand_offset);
+
+      tree address_base = build_address_operand (*brig_inst, *addr_operand);
+
+      uint32_t address_offset = 0;
+      while (bytes > 0)
+	{
+	  BrigOperandOffset32_t offset = *operand_ptr;
+	  const BrigBase *operand_element
+	    = m_parent.get_brig_operand_entry (offset);
+	  tree data
+	    = build_tree_operand (*brig_inst, *operand_element, instr_type);
+
+	  tree ptr_offset = build_int_cst (size_type_node, address_offset);
+	  tree address = build2 (POINTER_PLUS_EXPR, TREE_TYPE (address_base),
+				 address_base, ptr_offset);
+
+	  if (is_store && TREE_TYPE (data) != instr_type)
+	    {
+	      if (int_size_in_bytes (TREE_TYPE (data))
+		    == int_size_in_bytes (instr_type)
+		  && !INTEGRAL_TYPE_P (instr_type))
+		data = build1 (VIEW_CONVERT_EXPR, instr_type, data);
+	      else
+		data = convert (instr_type, data);
+	    }
+
+	  build_mem_access (brig_inst, address, data);
+
+	  address_offset += int_size_in_bytes (instr_type);
+	  ++operand_ptr;
+	  bytes -= 4;
+	}
+    }
+  else
+    {
+      tree_stl_vec operands = build_operands (*brig_inst);
+
+      tree &data = operands.at (0);
+      tree &addr = operands.at (1);
+      build_mem_access (brig_inst, addr, data);
+    }
+
+  return base->byteCount;
+}

gcc/brig/brigfrontend/brig-module-handler.cc

+/* brig-module-handler.cc -- brig module directive handling
+   Copyright (C) 2016 Free Software Foundation, Inc.
+   Contributed by Pekka Jaaskelainen <pekka.jaaskelainen@parmance.com>
+   for General Processor Tech.
+
+   This file is part of GCC.
+
+   GCC is free software; you can redistribute it and/or modify it under
+   the terms of the GNU General Public License as published by the Free
+   Software Foundation; either version 3, or (at your option) any later
+   version.
+
+   GCC is distributed in the hope that it will be useful, but WITHOUT ANY
+   WARRANTY; without even the implied warranty of MERCHANTABILITY or
+   FITNESS FOR A PARTICULAR PURPOSE.  See the GNU General Public License
+   for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with GCC; see the file COPYING3.  If not see
+   <http://www.gnu.org/licenses/>.  */
+
+#include "brig-code-entry-handler.h"
+#include "diagnostic-core.h"
+
+size_t
+brig_directive_module_handler::operator () (const BrigBase *base)
+{
+  const BrigDirectiveModule* mod = (const BrigDirectiveModule*)base;
+  m_parent.m_module_name = m_parent.get_string (mod->name).substr (1);
+  if (mod->hsailMajor != 1 || mod->hsailMinor != 0)
+    fatal_error (UNKNOWN_LOCATION, PHSA_ERROR_PREFIX_INCOMPATIBLE_MODULE " "
+		 "HSAIL version not supported. HSAIL 1.0 required.");
+  if (mod->machineModel != BRIG_MACHINE_LARGE)
+    fatal_error (UNKNOWN_LOCATION, PHSA_ERROR_PREFIX_INCOMPATIBLE_MODULE " "
+		 "Only HSA 'large' machine model supported.");
+  /* Do not check for the profile as the runtime conformance suite tests
+     with 'full' profile BRIGs even though they don't use any full profile
+     features.  This allows us to run the conformance suite with the
+     BRIG FE.  */
+  return base->byteCount;
+}

gcc/brig/brigfrontend/brig-queue-inst-handler.cc

+/* brig-queue-inst-handler.cc -- brig user mode queue related instruction
+   handling
+   Copyright (C) 2016 Free Software Foundation, Inc.
+   Contributed by Pekka Jaaskelainen <pekka.jaaskelainen@parmance.com>
+   for General Processor Tech.
+
+This file is part of GCC.
+
+GCC is free software; you can redistribute it and/or modify it under
+the terms of the GNU General Public License as published by the Free
+Software Foundation; either version 3, or (at your option) any later
+version.
+
+GCC is distributed in the hope that it will be useful, but WITHOUT ANY
+WARRANTY; without even the implied warranty of MERCHANTABILITY or
+FITNESS FOR A PARTICULAR PURPOSE.  See the GNU General Public License
+for more details.
+
+You should have received a copy of the GNU General Public License
+along with GCC; see the file COPYING3.  If not see
+<http://www.gnu.org/licenses/>.  */
+
+#include <sstream>
+
+#include "brig-code-entry-handler.h"
+#include "brig-util.h"
+#include "convert.h"
+#include "tree-pretty-print.h"
+#include "errors.h"
+#include "diagnostic-core.h"
+#include "brig-builtins.h"
+
+brig_queue_inst_handler::brig_queue_inst_handler (brig_to_generic &parent)
+  : brig_code_entry_handler (parent)
+{
+}
+
+size_t
+brig_queue_inst_handler::operator () (const BrigBase *base)
+{
+  const BrigInstBase &inst_base = *(const BrigInstBase *) base;
+
+  tree_stl_vec operands = build_operands (inst_base);
+
+  if (inst_base.opcode == BRIG_OPCODE_LDQUEUEWRITEINDEX
+      || inst_base.opcode == BRIG_OPCODE_LDQUEUEREADINDEX)
+    {
+      tree builtin
+	= inst_base.opcode == BRIG_OPCODE_LDQUEUEWRITEINDEX
+	? builtin_decl_explicit (BUILT_IN_HSAIL_LDQUEUEWRITEINDEX)
+	: builtin_decl_explicit (BUILT_IN_HSAIL_LDQUEUEREADINDEX);
+
+      tree expr
+	= call_builtin (builtin, 1, uint64_type_node,
+			uint64_type_node, operands[1]);
+      build_output_assignment (inst_base, operands[0], expr);
+    }
+  else if (inst_base.opcode == BRIG_OPCODE_STQUEUEWRITEINDEX
+	   || inst_base.opcode == BRIG_OPCODE_STQUEUEREADINDEX)
+    {
+      tree builtin
+	= inst_base.opcode == BRIG_OPCODE_STQUEUEWRITEINDEX
+	? builtin_decl_explicit (BUILT_IN_HSAIL_STQUEUEWRITEINDEX)
+	: builtin_decl_explicit (BUILT_IN_HSAIL_STQUEUEREADINDEX);
+
+      call_builtin (builtin, 2, void_type_node,
+		    uint64_type_node, operands[0], uint64_type_node,
+		    operands[1]);
+    }
+  else if (inst_base.opcode == BRIG_OPCODE_ADDQUEUEWRITEINDEX)
+    {
+      tree builtin = builtin_decl_explicit (BUILT_IN_HSAIL_ADDQUEUEWRITEINDEX);
+
+      tree expr = call_builtin (builtin, 2,
+				uint64_type_node, uint64_type_node, operands[1],
+				uint64_type_node, operands[2]);
+      build_output_assignment (inst_base, operands[0], expr);
+    }
+  else if (inst_base.opcode == BRIG_OPCODE_CASQUEUEWRITEINDEX)
+    {
+      tree builtin = builtin_decl_explicit (BUILT_IN_HSAIL_CASQUEUEWRITEINDEX);
+
+      tree expr
+	= call_builtin (builtin, 3, uint64_type_node,
+			uint64_type_node, operands[1], uint64_type_node,
+			operands[2], uint64_type_node, operands[3]);
+      build_output_assignment (inst_base, operands[0], expr);
+    }
+  else
+    gcc_unreachable ();
+
+  return base->byteCount;
+}

gcc/brig/brigfrontend/brig-seg-inst-handler.cc

+/* brig-seg-inst-handler.cc -- brig segment related instruction handling
+   Copyright (C) 2016 Free Software Foundation, Inc.
+   Contributed by Pekka Jaaskelainen <pekka.jaaskelainen@parmance.com>
+   for General Processor Tech.
+
+This file is part of GCC.
+
+GCC is free software; you can redistribute it and/or modify it under
+the terms of the GNU General Public License as published by the Free
+Software Foundation; either version 3, or (at your option) any later
+version.
+
+GCC is distributed in the hope that it will be useful, but WITHOUT ANY
+WARRANTY; without even the implied warranty of MERCHANTABILITY or
+FITNESS FOR A PARTICULAR PURPOSE.  See the GNU General Public License
+for more details.
+
+You should have received a copy of the GNU General Public License
+along with GCC; see the file COPYING3.  If not see
+<http://www.gnu.org/licenses/>.  */
+
+#include <sstream>
+
+#include "brig-code-entry-handler.h"
+#include "brig-util.h"
+#include "convert.h"
+#include "tree-pretty-print.h"
+#include "errors.h"
+#include "diagnostic-core.h"
+
+brig_seg_inst_handler::brig_seg_inst_handler (brig_to_generic &parent)
+  : brig_code_entry_handler (parent)
+{
+}
+
+size_t
+brig_seg_inst_handler::operator () (const BrigBase *base)
+{
+  const BrigInstBase &inst_base = *(const BrigInstBase *) base;
+
+  std::vector<tree> operands = build_operands (inst_base);
+
+  tree expr = NULL_TREE;
+
+  if (inst_base.opcode == BRIG_OPCODE_STOF)
+    {
+      const BrigInstSegCvt &inst = *(const BrigInstSegCvt *) base;
+
+      if (inst.segment == BRIG_SEGMENT_GROUP)
+	expr = build2 (PLUS_EXPR, size_type_node,
+		       convert_to_integer (size_type_node,
+					   m_parent.m_cf->m_group_base_arg),
+		       convert_to_integer (size_type_node, operands[1]));
+      else if (inst.segment == BRIG_SEGMENT_PRIVATE
+	       || inst.segment == BRIG_SEGMENT_SPILL)
+	expr = build2 (PLUS_EXPR, size_type_node,
+		       convert_to_integer (size_type_node,
+					   m_parent.m_cf->m_private_base_arg),
+		       convert_to_integer (size_type_node, operands[1]));
+      else
+       gcc_unreachable ();
+
+      if (!(inst.modifier & BRIG_SEG_CVT_NONULL))
+	{
+	  /* Need to convert the null value. -1 is used for 32b segments,
+	     and 0 for flat/global.  */
+	  tree cmp
+	    = build2 (EQ_EXPR, uint32_type_node,
+		      build_int_cstu (uint32_type_node, -1), operands[1]);
+
+	  tree null_check = build3 (COND_EXPR, size_type_node, cmp,
+				    build_int_cstu (size_type_node, 0), expr);
+
+	  expr = null_check;
+	}
+    }
+  else if (inst_base.opcode == BRIG_OPCODE_FTOS)
+    {
+      const BrigInstSegCvt &inst = *(const BrigInstSegCvt *) base;
+
+      if (inst.segment == BRIG_SEGMENT_GROUP)
+	expr = build2 (MINUS_EXPR, size_type_node,
+		       convert_to_integer (size_type_node,
+					   m_parent.m_cf->m_group_base_arg),
+		       convert_to_integer (size_type_node, operands[1]));
+      else if (inst.segment == BRIG_SEGMENT_PRIVATE)
+	expr = build2 (MINUS_EXPR, size_type_node,
+		       convert_to_integer (size_type_node,
+					   m_parent.m_cf->m_private_base_arg),
+		       convert_to_integer (size_type_node, operands[1]));
+      else
+	gcc_unreachable ();
+
+      if (!(inst.modifier & BRIG_SEG_CVT_NONULL))
+	{
+	  /* Need to convert the null value. -1 is used for 32b segments,
+	     and 0 for flat/global.  */
+	  tree cmp = build2 (EQ_EXPR, size_type_node,
+			     build_int_cstu (size_type_node, 0), operands[1]);
+
+	  tree null_check
+	    = build3 (COND_EXPR, size_type_node, cmp,
+		      build_int_cstu (uint32_type_node, -1), expr);
+	  expr = null_check;
+	}
+    }
+  else if (inst_base.opcode == BRIG_OPCODE_NULLPTR)
+    {
+      const BrigInstSeg &inst = *(const BrigInstSeg *) base;
+      if (inst.segment == BRIG_SEGMENT_GLOBAL
+	  || inst.segment == BRIG_SEGMENT_FLAT
+	  || inst.segment == BRIG_SEGMENT_READONLY)
+	expr = build_int_cstu (uint64_type_node, 0);
+      else
+	expr = build_int_cstu (uint32_type_node, -1);
+    }
+  else if (inst_base.opcode == BRIG_OPCODE_SEGMENTP)
+    {
+      const BrigInstSegCvt &inst = *(const BrigInstSegCvt *) base;
+
+      tree builtin = NULL_TREE;
+      switch (inst.segment)
+	{
+	case BRIG_SEGMENT_GLOBAL:
+	  builtin = builtin_decl_explicit (BUILT_IN_HSAIL_SEGMENTP_GLOBAL);
+	  break;
+	case BRIG_SEGMENT_GROUP:
+	  builtin = builtin_decl_explicit (BUILT_IN_HSAIL_SEGMENTP_GROUP);
+	  break;
+	case BRIG_SEGMENT_PRIVATE:
+	  builtin = builtin_decl_explicit (BUILT_IN_HSAIL_SEGMENTP_PRIVATE);
+	  break;
+	default:
+	  gcc_unreachable ();
+	}
+
+      expr = call_builtin (builtin, 2,
+			   uint32_type_node, uint64_type_node, operands[1],
+			   ptr_type_node, m_parent.m_cf->m_context_arg);
+    }
+  else
+    gcc_unreachable ();
+
+  build_output_assignment (inst_base, operands[0], expr);
+  return base->byteCount;
+}

gcc/brig/brigfrontend/brig-signal-inst-handler.cc

+/* brig-signal-inst-handler.cc -- brig signal instruction handling
+   Copyright (C) 2016 Free Software Foundation, Inc.
+   Contributed by Pekka Jaaskelainen <pekka.jaaskelainen@parmance.com>
+   for General Processor Tech.
+
+This file is part of GCC.
+
+GCC is free software; you can redistribute it and/or modify it under
+the terms of the GNU General Public License as published by the Free
+Software Foundation; either version 3, or (at your option) any later
+version.
+
+GCC is distributed in the hope that it will be useful, but WITHOUT ANY
+WARRANTY; without even the implied warranty of MERCHANTABILITY or
+FITNESS FOR A PARTICULAR PURPOSE.  See the GNU General Public License
+for more details.
+
+You should have received a copy of the GNU General Public License
+along with GCC; see the file COPYING3.  If not see
+<http://www.gnu.org/licenses/>.  */
+
+#include <sstream>
+
+#include "brig-code-entry-handler.h"
+#include "brig-util.h"
+#include "fold-const.h"
+#include "diagnostic.h"
+#include "tree-pretty-print.h"
+#include "print-tree.h"
+#include "convert.h"
+#include "langhooks.h"
+#include "gimple-expr.h"
+
+size_t
+brig_signal_inst_handler::operator () (const BrigBase *base)
+{
+  const BrigInstSignal *inst = (const BrigInstSignal *) base;
+  BrigAtomicOperation8_t atomic_opcode;
+  atomic_opcode = inst->signalOperation;
+
+  return generate_tree (inst->base, atomic_opcode);
+}

gcc/brig/brigfrontend/brig-to-generic.h

+/* brig-to-generic.h -- brig to gcc generic conversion
+   Copyright (C) 2016 Free Software Foundation, Inc.
+   Contributed by Pekka Jaaskelainen <pekka.jaaskelainen@parmance.com>
+   for General Processor Tech.
+
+   This file is part of GCC.
+
+   GCC is free software; you can redistribute it and/or modify it under
+   the terms of the GNU General Public License as published by the Free
+   Software Foundation; either version 3, or (at your option) any later
+   version.
+
+   GCC is distributed in the hope that it will be useful, but WITHOUT ANY
+   WARRANTY; without even the implied warranty of MERCHANTABILITY or
+   FITNESS FOR A PARTICULAR PURPOSE.  See the GNU General Public License
+   for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with GCC; see the file COPYING3.  If not see
+   <http://www.gnu.org/licenses/>.  */
+
+#ifndef BRIG_TO_GENERIC_H
+#define BRIG_TO_GENERIC_H
+
+#include <string>
+#include <map>
+#include <vector>
+
+#include "config.h"
+#include "system.h"
+#include "ansidecl.h"
+#include "coretypes.h"
+#include "opts.h"
+#include "tree.h"
+#include "tree-iterator.h"
+#include "hsa-brig-format.h"
+#include "brig-function.h"
+
+
+struct reg_decl_index_entry;
+
+/* Converts an HSAIL BRIG input to GENERIC.  This class holds global state
+   for the translation process.  Handling of the smaller pieces of BRIG data
+   is delegated to various handler classes declared in
+   brig-code-entry-handlers.h.  */
+
+class brig_to_generic
+{
+public:
+  typedef std::map<const BrigDirectiveVariable *, tree> variable_index;
+
+private:
+  typedef std::map<std::string, size_t> var_offset_table;
+  typedef std::map<const BrigBase *, std::string> name_index;
+
+public:
+  brig_to_generic ();
+
+  void parse (const char *brig_blob);
+
+  void write_globals ();
+
+  std::string get_string (size_t entry_offset) const;
+
+  const BrigData *get_brig_data_entry (size_t entry_offset) const;
+  const BrigBase *get_brig_operand_entry (size_t entry_offset) const;
+  const BrigBase *get_brig_code_entry (size_t entry_offset) const;
+
+  void append_global (tree g);
+
+  tree function_decl (const std::string &name);
+  void add_function_decl (const std::string &name, tree func_decl);
+
+  tree global_variable (const std::string &name) const;
+  void add_global_variable (const std::string &name, tree var_decl);
+  void add_host_def_var_ptr (const std::string &name, tree var_decl);
+
+  void start_function (tree f);
+  void finish_function ();
+
+  void append_group_variable (const std::string &name, size_t size,
+			      size_t alignment);
+
+  void append_private_variable (const std::string &name, size_t size,
+				size_t alignment);
+
+  size_t group_variable_segment_offset (const std::string &name) const;
+
+  bool
+  has_group_variable (const std::string &name) const;
+
+  size_t
+  private_variable_segment_offset (const std::string &name) const;
+
+  bool
+  has_private_variable (const std::string &name) const;
+
+  size_t private_variable_size (const std::string &name) const;
+
+  template <typename T>
+    std::string
+    get_mangled_name_tmpl (const T *brigVar) const;
+
+  std::string get_mangled_name (const BrigDirectiveFbarrier *fbar) const
+    { return get_mangled_name_tmpl (fbar); }
+  std::string get_mangled_name (const BrigDirectiveVariable *var) const
+    { return get_mangled_name_tmpl (var); }
+  std::string get_mangled_name (const BrigDirectiveExecutable *func) const;
+
+  size_t group_segment_size () const;
+  size_t private_segment_size () const;
+
+  brig_function *get_finished_function (tree func_decl);
+
+  static tree s_fp16_type;
+  static tree s_fp32_type;
+  static tree s_fp64_type;
+
+  /* The default rounding mode that should be used for float instructions.
+     This can be set in each BRIG module header.  */
+  BrigRound8_t m_default_float_rounding_mode;
+
+  /* The currently built function.  */
+  brig_function *m_cf;
+
+  /* The name of the currently handled BRIG module.  */
+  std::string m_module_name;
+
+private:
+  /* The BRIG blob and its different sections of the file currently being
+     parsed.  */
+  const char *m_brig;
+  const char *m_data;
+  size_t m_data_size;
+  const char *m_operand;
+  size_t m_operand_size;
+  const char *m_code;
+  size_t m_code_size;
+
+  tree m_globals;
+
+  label_index m_global_variables;
+
+  /* The size of each private variable, including the alignment padding.  */
+  std::map<std::string, size_t> m_private_data_sizes;
+
+  /* The same for group variables.  */
+  size_t m_next_group_offset;
+  var_offset_table m_group_offsets;
+
+  /* And private.  */
+  size_t m_next_private_offset;
+  var_offset_table m_private_offsets;
+
+  /* Name index for declared functions.  */
+  label_index m_function_index;
+
+  /* Stores all processed kernels in order.  */
+  std::vector<brig_function *> m_kernels;
+
+  /* Stores all already processed functions from the translation unit
+     for some interprocedural analysis.  */
+  std::map<std::string, brig_function *> m_finished_functions;
+
+  /* The parsed BRIG blobs.  Owned and will be deleted after use.  */
+  std::vector<const char *> m_brig_blobs;
+
+  /* The original dump file.  */
+  FILE *m_dump_file;
+
+  /* The original dump file flags.  */
+  int m_dump_flags;
+};
+
+/* Produce a "mangled name" for the given brig variable.  The mangling is used
+   to make unique global symbol names for module and function scope variables.
+   The templated version is suitable for most of the variable types.  Functions
+   and kernels (BrigDirectiveExecutable) are handled with a specialized
+   get_mangled_name() version.  */
+
+template <typename T>
+std::string
+brig_to_generic::get_mangled_name_tmpl (const T *brigVar) const
+{
+  std::string var_name = get_string (brigVar->name).substr (1);
+
+  /* Mangle the variable name using the function name and the module name
+     in case of a function scope variable.  */
+  if (m_cf != NULL
+      && m_cf->has_function_scope_var (&brigVar->base))
+    var_name = m_cf->m_name + "." + var_name;
+
+  if (brigVar->linkage == BRIG_LINKAGE_MODULE)
+    var_name = "gccbrig." + m_module_name + "." + var_name;
+  return var_name;
+}
+
+/* An interface to organize the different types of BRIG element handlers.  */
+
+class brig_entry_handler
+{
+public:
+  brig_entry_handler (brig_to_generic &parent) : m_parent (parent)
+  {
+  }
+
+  /* Handles the brig_code data at the given pointer and adds it to the
+     currently built tree.  Returns the number of consumed bytes;  */
+  virtual size_t operator () (const BrigBase *base) = 0;
+
+protected:
+  brig_to_generic &m_parent;
+};
+
+tree call_builtin (tree pdecl, int nargs, tree rettype, ...);
+
+tree build_reinterpret_cast (tree destination_type, tree source);
+
+tree build_stmt (enum tree_code code, ...);
+
+tree get_unsigned_int_type (tree type);
+
+void dump_function (FILE *dump_file, brig_function *f);
+
+#endif

gcc/brig/brigfrontend/brig-util.h

+/* brig-util.h -- gccbrig utility functions
+   Copyright (C) 2016 Free Software Foundation, Inc.
+   Contributed by Pekka Jaaskelainen <pekka.jaaskelainen@parmance.com>
+   for General Processor Tech.
+
+This file is part of GCC.
+
+GCC is free software; you can redistribute it and/or modify it under
+the terms of the GNU General Public License as published by the Free
+Software Foundation; either version 3, or (at your option) any later
+version.
+
+GCC is distributed in the hope that it will be useful, but WITHOUT ANY
+WARRANTY; without even the implied warranty of MERCHANTABILITY or
+FITNESS FOR A PARTICULAR PURPOSE.  See the GNU General Public License
+for more details.
+
+You should have received a copy of the GNU General Public License
+along with GCC; see the file COPYING3.  If not see
+<http://www.gnu.org/licenses/>.  */
+
+#ifndef GCC_BRIG_UTIL_H
+#define GCC_BRIG_UTIL_H
+
+#include "brig-to-generic.h"
+
+bool gccbrig_hsa_opcode_op_output_p (BrigOpcode16_t opcode, int opnum);
+
+unsigned gccbrig_hsa_type_bit_size (BrigType16_t t);
+
+uint64_t gccbrig_to_uint64_t (const BrigUInt64 &brig_type);
+
+int gccbrig_reg_size (const BrigOperandRegister *brig_reg);
+
+std::string gccbrig_reg_name (const BrigOperandRegister *reg);
+
+std::string gccbrig_type_name (BrigType16_t type);
+
+std::string gccbrig_segment_name (BrigSegment8_t segment);
+
+bool gccbrig_is_float_type (BrigType16_t type);
+
+bool gccbrig_is_bit_operation (BrigOpcode16_t opcode);
+
+BrigType16_t gccbrig_tree_type_to_hsa_type (tree tree_type);
+tree gccbrig_tree_type_for_hsa_type (BrigType16_t brig_type);
+
+bool gccbrig_might_be_host_defined_var_p (const BrigDirectiveVariable *brigVar);
+
+/* From hsa.h.  */
+bool hsa_type_packed_p (BrigType16_t type);
+
+#endif

gcc/brig/brigfrontend/brig-variable-handler.cc

+/* brig-variable-handler.cc -- brig variable directive handling
+   Copyright (C) 2016 Free Software Foundation, Inc.
+   Contributed by Pekka Jaaskelainen <pekka.jaaskelainen@parmance.com>
+   for General Processor Tech.
+
+   This file is part of GCC.
+
+   GCC is free software; you can redistribute it and/or modify it under
+   the terms of the GNU General Public License as published by the Free
+   Software Foundation; either version 3, or (at your option) any later
+   version.
+
+   GCC is distributed in the hope that it will be useful, but WITHOUT ANY
+   WARRANTY; without even the implied warranty of MERCHANTABILITY or
+   FITNESS FOR A PARTICULAR PURPOSE.  See the GNU General Public License
+   for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with GCC; see the file COPYING3.  If not see
+   <http://www.gnu.org/licenses/>.  */
+
+#include "brig-code-entry-handler.h"
+
+#include "stringpool.h"
+#include "errors.h"
+#include "brig-machine.h"
+#include "brig-util.h"
+#include "print-tree.h"
+#include "diagnostic-core.h"
+
+tree
+brig_directive_variable_handler::build_variable
+  (const BrigDirectiveVariable *brigVar, tree_code var_decltype)
+{
+  std::string var_name = m_parent.get_mangled_name (brigVar);
+
+  bool is_definition = brigVar->modifier & BRIG_VARIABLE_DEFINITION;
+
+  tree name_identifier = get_identifier (var_name.c_str ());
+
+  tree var_decl;
+  tree t;
+  if (brigVar->type & BRIG_TYPE_ARRAY)
+    {
+      tree element_type
+	= gccbrig_tree_type_for_hsa_type (brigVar->type & ~BRIG_TYPE_ARRAY);
+      uint64_t element_count = gccbrig_to_uint64_t (brigVar->dim);
+      if (is_definition && element_count == 0)
+	fatal_error (UNKNOWN_LOCATION, "Array definition with zero elements.");
+      if (var_decltype == PARM_DECL)
+	t = build_pointer_type (element_type);
+      else
+	t = build_array_type_nelts (element_type, element_count);
+    }
+  else
+    {
+      t = gccbrig_tree_type_for_hsa_type (brigVar->type);
+    }
+
+  size_t alignment = get_brig_var_alignment (brigVar);
+
+  if (brigVar->segment == BRIG_SEGMENT_READONLY
+      || brigVar->segment == BRIG_SEGMENT_KERNARG
+      || (brigVar->modifier & BRIG_VARIABLE_CONST))
+    TYPE_READONLY (t) = 1;
+
+  TYPE_ADDR_SPACE (t) = gccbrig_get_target_addr_space_id (brigVar->segment);
+
+  var_decl = build_decl (UNKNOWN_LOCATION, var_decltype, name_identifier, t);
+
+  SET_DECL_ALIGN (var_decl, alignment * BITS_PER_UNIT);
+
+  /* Force the HSA alignments.  */
+  DECL_USER_ALIGN (var_decl) = 1;
+
+  TREE_USED (var_decl) = 1;
+
+  TREE_PUBLIC (var_decl) = 1;
+  if (is_definition)
+    DECL_EXTERNAL (var_decl) = 0;
+  else
+    DECL_EXTERNAL (var_decl) = 1; /* The definition is elsewhere.  */
+
+  if (brigVar->init != 0)
+    {
+      gcc_assert (brigVar->segment == BRIG_SEGMENT_READONLY
+		  || brigVar->segment == BRIG_SEGMENT_GLOBAL);
+
+      const BrigBase *cst_operand_data
+	= m_parent.get_brig_operand_entry (brigVar->init);
+
+      tree initializer = NULL_TREE;
+      if (cst_operand_data->kind == BRIG_KIND_OPERAND_CONSTANT_BYTES)
+	initializer = get_tree_cst_for_hsa_operand
+	  ((const BrigOperandConstantBytes *) cst_operand_data, t);
+      else
+	error ("variable initializers of type %x not implemented",
+	       cst_operand_data->kind);
+      gcc_assert (initializer != NULL_TREE);
+      DECL_INITIAL (var_decl) = initializer;
+    }
+
+  if (var_decltype == PARM_DECL)
+    {
+      DECL_ARG_TYPE (var_decl) = TREE_TYPE (var_decl);
+      DECL_EXTERNAL (var_decl) = 0;
+      TREE_PUBLIC (var_decl) = 0;
+    }
+
+  TREE_ADDRESSABLE (var_decl) = 1;
+
+  TREE_USED (var_decl) = 1;
+  DECL_NONLOCAL (var_decl) = 1;
+  DECL_ARTIFICIAL (var_decl) = 0;
+
+  return var_decl;
+}
+
+size_t
+brig_directive_variable_handler::operator () (const BrigBase *base)
+{
+  const BrigDirectiveVariable *brigVar = (const BrigDirectiveVariable *) base;
+
+  bool is_definition = brigVar->modifier & BRIG_VARIABLE_DEFINITION;
+
+  size_t var_size;
+  tree var_type;
+  if (brigVar->type & BRIG_TYPE_ARRAY)
+    {
+      tree element_type
+	= gccbrig_tree_type_for_hsa_type (brigVar->type & ~BRIG_TYPE_ARRAY);
+      uint64_t element_count = gccbrig_to_uint64_t (brigVar->dim);
+      if (is_definition && element_count == 0)
+	fatal_error (UNKNOWN_LOCATION, "Array definition with zero elements.");
+      var_type = build_array_type_nelts (element_type, element_count);
+      size_t element_size = tree_to_uhwi (TYPE_SIZE (element_type));
+      var_size = element_size * element_count / 8;
+    }
+  else
+    {
+      var_type = gccbrig_tree_type_for_hsa_type (brigVar->type);
+      var_size = tree_to_uhwi (TYPE_SIZE (var_type)) / 8;
+    }
+
+  size_t alignment = get_brig_var_alignment (brigVar);
+
+  if (m_parent.m_cf != NULL)
+    m_parent.m_cf->m_function_scope_vars.insert (base);
+
+  std::string var_name = m_parent.get_mangled_name (brigVar);
+  if (brigVar->segment == BRIG_SEGMENT_KERNARG)
+    {
+      /* Do not create a real variable, but only a table of
+	 offsets to the kernarg segment buffer passed as the
+	 single argument by the kernel launcher for later
+	 reference.  Ignore kernel declarations.  */
+      if (m_parent.m_cf != NULL && m_parent.m_cf->m_func_decl != NULL_TREE)
+	m_parent.m_cf->append_kernel_arg (brigVar, var_size, alignment);
+      return base->byteCount;
+    }
+  else if (brigVar->segment == BRIG_SEGMENT_GROUP)
+    {
+      /* Handle group region variables similarly as kernargs:
+	 assign offsets to the group region on the fly when
+	 a new module scope or function scope group variable is
+	 introduced.  These offsets will be then added to the
+	 group_base hidden pointer passed to the kernel in order to
+	 get the flat address.  */
+      if (!m_parent.has_group_variable (var_name))
+	m_parent.append_group_variable (var_name, var_size, alignment);
+      return base->byteCount;
+    }
+  else if (brigVar->segment == BRIG_SEGMENT_PRIVATE
+	   || brigVar->segment == BRIG_SEGMENT_SPILL)
+    {
+      /* Private variables are handled like group variables,
+	 except that their offsets are multiplied by the work-item
+	 flat id, when accessed.  */
+      if (!m_parent.has_private_variable (var_name))
+	m_parent.append_private_variable (var_name, var_size, alignment);
+      return base->byteCount;
+    }
+  else if (brigVar->segment == BRIG_SEGMENT_GLOBAL
+	   || brigVar->segment == BRIG_SEGMENT_READONLY)
+    {
+      tree def = is_definition ? NULL_TREE :
+	m_parent.global_variable (var_name);
+
+      if (!is_definition && def != NULL_TREE)
+	{
+	  /* We have a definition already for this declaration.
+	     Use the definition instead of the declaration.  */
+	}
+      else if (gccbrig_might_be_host_defined_var_p (brigVar))
+	{
+	  tree var_decl = build_variable (brigVar);
+	  m_parent.add_host_def_var_ptr (var_name, var_decl);
+	}
+      else
+	{
+	  tree var_decl = build_variable (brigVar);
+	  /* Make all global variables program scope for now
+	     so we can get their address from the Runtime API.  */
+	  DECL_CONTEXT (var_decl) = NULL_TREE;
+	  TREE_STATIC (var_decl) = 1;
+	  m_parent.add_global_variable (var_name, var_decl);
+	}
+    }
+  else if (brigVar->segment == BRIG_SEGMENT_ARG)
+    {
+
+      if (m_parent.m_cf->m_generating_arg_block)
+	{
+	  tree var_decl = build_variable (brigVar);
+	  tree bind_expr = m_parent.m_cf->m_current_bind_expr;
+
+	  DECL_CONTEXT (var_decl) = m_parent.m_cf->m_func_decl;
+	  DECL_CHAIN (var_decl) = BIND_EXPR_VARS (bind_expr);
+	  BIND_EXPR_VARS (bind_expr) = var_decl;
+	  TREE_PUBLIC (var_decl) = 0;
+
+	  m_parent.m_cf->add_arg_variable (brigVar, var_decl);
+	}
+      else
+	{
+	  /* Must be an incoming function argument which has
+	     been parsed in brig-function-handler.cc.  No
+	     need to generate anything here.  */
+	}
+    }
+  else
+    gcc_unreachable ();
+
+  return base->byteCount;
+}
+
+/* Returns the alignment for the given BRIG variable.  In case the variable
+   explicitly defines alignment and its larger than the natural alignment,
+   returns it instead of the natural one.  */
+
+size_t
+brig_directive_variable_handler::get_brig_var_alignment
+(const BrigDirectiveVariable *brigVar)
+{
+
+  size_t defined_alignment
+    = brigVar->align == BRIG_ALIGNMENT_NONE ? 0 : 1 << (brigVar->align - 1);
+  size_t natural_alignment;
+  if (brigVar->type & BRIG_TYPE_ARRAY)
+    {
+      tree element_type
+	= gccbrig_tree_type_for_hsa_type (brigVar->type & ~BRIG_TYPE_ARRAY);
+      size_t element_size = tree_to_uhwi (TYPE_SIZE (element_type));
+      natural_alignment = element_size / BITS_PER_UNIT;
+    }
+  else
+    {
+      tree t = gccbrig_tree_type_for_hsa_type (brigVar->type);
+      natural_alignment = tree_to_uhwi (TYPE_SIZE (t)) / BITS_PER_UNIT;
+    }
+
+  return natural_alignment > defined_alignment
+    ? natural_alignment : defined_alignment;
+}

gcc/brig/brigfrontend/phsa.h

+/* phsa.h -- interfacing between the gcc BRIG FE and the phsa runtime
+   Copyright (C) 2016 Free Software Foundation, Inc.
+   Contributed by Pekka Jaaskelainen <pekka.jaaskelainen@parmance.com>
+   for General Processor Tech.
+
+This file is part of GCC.
+
+GCC is free software; you can redistribute it and/or modify it under
+the terms of the GNU General Public License as published by the Free
+Software Foundation; either version 3, or (at your option) any later
+version.
+
+GCC is distributed in the hope that it will be useful, but WITHOUT ANY
+WARRANTY; without even the implied warranty of MERCHANTABILITY or
+FITNESS FOR A PARTICULAR PURPOSE.  See the GNU General Public License
+for more details.
+
+You should have received a copy of the GNU General Public License
+along with GCC; see the file COPYING3.  If not see
+<http://www.gnu.org/licenses/>.  */
+
+#ifndef PHSA_H
+#define PHSA_H
+
+#include <stdint.h>
+
+/* This struct is used to pass information from the BRIG FE to the
+   runtime of the finalizer kernel, its control directives etc.
+   The data is passed raw in a special ELF section named
+   phsa.kerneldesc.kernel_function_name.  */
+
+typedef struct __attribute__((__packed__))
+{
+  /* Set to 1 in case the function is a kernel.  */
+  uint8_t is_kernel;
+  /* The size of the group segment used by the kernel.  */
+  uint32_t group_segment_size;
+  /* Size of the private segment used by a single work-item.  */
+  uint32_t private_segment_size;
+  /* Total size of the kernel arguments.  */
+  uint32_t kernarg_segment_size;
+  /* Maximum alignment of a kernel argument variable.  */
+  uint16_t kernarg_max_align;
+  /* Maximum size (in bytes) of dynamic group memory.  */
+  uint32_t max_dynamic_group_size;
+  /* Max number of work-items used to launch the kernel.  */
+  uint64_t max_flat_grid_size;
+  /* Max number of work-items in a work-group used to launch the kernel.  */
+  uint32_t max_flat_workgroup_size;
+  /* The grid size required by the kernel.  */
+  uint64_t required_grid_size[3];
+  /* The work group size required by the kernel.  */
+  uint32_t required_workgroup_size[3];
+  /* The number of dimensions required by the kernel.  */
+  uint8_t required_dim;
+
+} phsa_descriptor;
+
+/* The prefix to use in the ELF section containing descriptor for
+   a function.  */
+#define PHSA_DESC_SECTION_PREFIX "phsa.desc."
+#define PHSA_HOST_DEF_PTR_PREFIX "__phsa.host_def."
+
+/* The frontend error messages are parsed by the host runtime, known
+   prefix strings are used to separate the different runtime error
+   codes.  */
+#define PHSA_ERROR_PREFIX_INCOMPATIBLE_MODULE "Incompatible module:"
+
+#endif

gcc/brig/brigspec.c

+/* brigspec.c -- Specific flags and argument handling of the gcc BRIG front end.
+   Copyright (C) 2016 Free Software Foundation, Inc.
+   Contributed by Pekka Jaaskelainen <pekka.jaaskelainen@parmance.com>
+   for General Processor Tech.
+
+This file is part of GCC.
+
+GCC is free software; you can redistribute it and/or modify it under
+the terms of the GNU General Public License as published by the Free
+Software Foundation; either version 3, or (at your option) any later
+version.
+
+GCC is distributed in the hope that it will be useful, but WITHOUT ANY
+WARRANTY; without even the implied warranty of MERCHANTABILITY or
+FITNESS FOR A PARTICULAR PURPOSE.  See the GNU General Public License
+for more details.
+
+You should have received a copy of the GNU General Public License
+along with GCC; see the file COPYING3.  If not see
+<http://www.gnu.org/licenses/>.  */
+
+#include "config.h"
+#include "system.h"
+#include "coretypes.h"
+#include "tm.h"
+#include "gcc.h"
+#include "opts.h"
+
+/* This bit is set if we saw a `-xfoo' language specification.  */
+#define LANGSPEC (1 << 1)
+/* This bit is set if they did `-lm' or `-lmath'.  */
+#define MATHLIB (1 << 2)
+/* This bit is set if they did `-lpthread'.  */
+#define THREADLIB (1 << 3)
+/* This bit is set if they did `-lc'.  */
+#define WITHLIBC (1 << 4)
+/* Skip this option.  */
+#define SKIPOPT (1 << 5)
+
+#ifndef MATH_LIBRARY
+#define MATH_LIBRARY "m"
+#endif
+#ifndef MATH_LIBRARY_PROFILE
+#define MATH_LIBRARY_PROFILE MATH_LIBRARY
+#endif
+
+#define LIBHSAIL "hsail-rt"
+
+void
+lang_specific_driver (struct cl_decoded_option **in_decoded_options,
+		      unsigned int *in_decoded_options_count,
+		      int *in_added_libraries)
+{
+  unsigned int i, j;
+
+  /* The new argument list will be contained in this.  */
+  struct cl_decoded_option *new_decoded_options;
+
+  /* An array used to flag each argument that needs a bit set for
+     LANGSPEC, MATHLIB, or WITHLIBC.  */
+  int *args;
+
+  /* By default, we throw on the math library if we have one.  */
+  int need_math = (MATH_LIBRARY[0] != '\0');
+
+  /* True if we should add -shared-libgcc to the command-line.  */
+  int shared_libgcc = 1;
+
+  /* The total number of arguments with the new stuff.  */
+  unsigned int argc;
+
+  /* The argument list.  */
+  struct cl_decoded_option *decoded_options;
+
+  /* The number of libraries added in.  */
+  int added_libraries;
+
+  /* The total number of arguments with the new stuff.  */
+  int num_args = 1;
+
+  argc = *in_decoded_options_count;
+  decoded_options = *in_decoded_options;
+  added_libraries = *in_added_libraries;
+
+  args = XCNEWVEC (int, argc);
+
+  for (i = 1; i < argc; i++)
+    {
+      switch (decoded_options[i].opt_index)
+	{
+	case OPT_o:
+	  break;
+
+	case OPT_SPECIAL_input_file:
+	  break;
+	}
+    }
+
+  /* Make sure to have room for the trailing NULL argument.  */
+  num_args = argc + need_math + shared_libgcc + 10;
+  new_decoded_options = XNEWVEC (struct cl_decoded_option, num_args);
+
+  i = 0;
+  j = 0;
+
+  /* Copy the 0th argument, i.e., the name of the program itself.  */
+  new_decoded_options[j++] = decoded_options[i++];
+
+  /* NOTE: We start at 1 now, not 0.  */
+  while (i < argc)
+    {
+      new_decoded_options[j] = decoded_options[i];
+
+      if ((args[i] & SKIPOPT) != 0)
+	--j;
+
+      i++;
+      j++;
+    }
+
+  generate_option (OPT_l, LIBHSAIL, 1, CL_DRIVER, &new_decoded_options[j]);
+  j++;
+
+  *in_decoded_options_count = j;
+  *in_decoded_options = new_decoded_options;
+  *in_added_libraries = added_libraries;
+}
+
+/* Called before linking.  Returns 0 on success and -1 on failure.  */
+
+int lang_specific_pre_link (void) /* Not used for Brig.  */ { return 0; }
+
+/* Number of extra output files that lang_specific_pre_link may generate.  */
+
+int lang_specific_extra_outfiles = 0; /* Not used for Brig.  */

gcc/brig/config-lang.in

+# config-lang.in -- Top level configure fragment for gcc BRIG (HSAIL) frontend.
+
+# Copyright (C) 2015 Free Software Foundation, Inc.
+
+# This file is part of GCC.
+
+# GCC is free software; you can redistribute it and/or modify
+# it under the terms of the GNU General Public License as published by
+# the Free Software Foundation; either version 3, or (at your option)
+# any later version.
+
+# GCC is distributed in the hope that it will be useful,
+# but WITHOUT ANY WARRANTY; without even the implied warranty of
+# MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+# GNU General Public License for more details.
+
+# You should have received a copy of the GNU General Public License
+# along with GCC; see the file COPYING3.  If not see
+# <http://www.gnu.org/licenses/>.
+
+# Configure looks for the existence of this file to auto-config each language.
+# We define several parameters used by configure:
+#
+# language	- name of language as it would appear in $(LANGUAGES)
+# compilers	- value to add to $(COMPILERS)
+
+language="brig"
+
+compilers="brig1\$(exeext)"
+
+target_libs="target-libbrig target-libhsail-rt"
+
+# The BRIG frontend is written in C++, so we need to build the C++
+# compiler during stage 1.  Note: when cross-compiling / not bootstrapping,
+# this can be safely removed. gcc 4.9.1 force enables c++/libstdc++ to the
+# target compiler due to this.
+lang_requires_boot_languages=c++
+
+gtfiles="\$(srcdir)/brig/brig-lang.c \$(srcdir)/brig/brig-c.h"
+
+build_by_default="no"

gcc/brig/lang-specs.h

+/* lang-specs.h -- gcc driver specs for BRIG (HSAIL) frontend.
+   Copyright (C) 2016 Free Software Foundation, Inc.
+   Contributed by Pekka Jaaskelainen <pekka.jaaskelainen@parmance.com>
+   for General Processor Tech.
+
+This file is part of GCC.
+
+GCC is free software; you can redistribute it and/or modify it under
+the terms of the GNU General Public License as published by the Free
+Software Foundation; either version 3, or (at your option) any later
+version.
+
+GCC is distributed in the hope that it will be useful, but WITHOUT ANY
+WARRANTY; without even the implied warranty of MERCHANTABILITY or
+FITNESS FOR A PARTICULAR PURPOSE.  See the GNU General Public License
+for more details.
+
+You should have received a copy of the GNU General Public License
+along with GCC; see the file COPYING3.  If not see
+<http://www.gnu.org/licenses/>.  */
+
+/* This is the contribution to the `default_compilers' array in gcc.c
+   for the BRIG (HSAIL) input.  */
+
+{".brig", "@brig", 0, 1, 0},
+  {"@brig",
+   "brig1 %i %(cc1_options) %{I*} %{L*} %D %{!fsyntax-only:%(invoke_as)}", 0, 1,
+   0},

gcc/brig/lang.opt

+; lang.opt -- Options for the gcc BRIG (HSAIL) front end.
+
+; Copyright (C) 2015 Free Software Foundation, Inc.
+;
+; This file is part of GCC.
+;
+; GCC is free software; you can redistribute it and/or modify it under
+; the terms of the GNU General Public License as published by the Free
+; Software Foundation; either version 3, or (at your option) any later
+; version.
+;
+; GCC is distributed in the hope that it will be useful, but WITHOUT ANY
+; WARRANTY; without even the implied warranty of MERCHANTABILITY or
+; FITNESS FOR A PARTICULAR PURPOSE.  See the GNU General Public License
+; for more details.
+;
+; You should have received a copy of the GNU General Public License
+; along with GCC; see the file COPYING3.  If not see
+; <http://www.gnu.org/licenses/>.
+
+; See the GCC internals manual for a description of this file's format.
+
+; Please try to keep this file in ASCII collating order.
+
+Language
+BRIG
+
+-dump
+BRIG Separate Alias(d)
+
+-dump=
+BRIG Joined Alias(d)
+
+L
+BRIG Joined Separate
+; Not documented
+
+-output=
+BRIG Driver Joined Alias(o) MissingArgError(missing filename after %qs)
+
+; This comment is to ensure we retain the blank line above.

gcc/builtin-types.def

@@ -67,7 +67,10 @@ DEF_PRIMITIVE_TYPE (BT_LONGLONG, long_long_integer_type_node)
 DEF_PRIMITIVE_TYPE (BT_ULONGLONG, long_long_unsigned_type_node)
 DEF_PRIMITIVE_TYPE (BT_INTMAX, intmax_type_node)
 DEF_PRIMITIVE_TYPE (BT_UINTMAX, uintmax_type_node)
-DEF_PRIMITIVE_TYPE (BT_UINT16, uint16_type_node)
+DEF_PRIMITIVE_TYPE (BT_INT8, signed_char_type_node)
+DEF_PRIMITIVE_TYPE (BT_INT16, short_integer_type_node)
+DEF_PRIMITIVE_TYPE (BT_UINT8, char_type_node)
+DEF_PRIMITIVE_TYPE (BT_UINT16, short_unsigned_type_node)
 DEF_PRIMITIVE_TYPE (BT_UINT32, uint32_type_node)
 DEF_PRIMITIVE_TYPE (BT_UINT64, uint64_type_node)
 DEF_PRIMITIVE_TYPE (BT_WORD, (*lang_hooks.types.type_for_mode) (word_mode, 1))
@@ -167,6 +170,7 @@ DEF_FUNCTION_TYPE_0 (BT_FN_CONST_STRING, BT_CONST_STRING)
 DEF_FUNCTION_TYPE_0 (BT_FN_PID, BT_PID)
 DEF_FUNCTION_TYPE_0 (BT_FN_INT, BT_INT)
 DEF_FUNCTION_TYPE_0 (BT_FN_UINT, BT_UINT)
+DEF_FUNCTION_TYPE_0 (BT_FN_ULONG, BT_ULONG)
 DEF_FUNCTION_TYPE_0 (BT_FN_FLOAT, BT_FLOAT)
 DEF_FUNCTION_TYPE_0 (BT_FN_DOUBLE, BT_DOUBLE)
 /* For "long double" we use LONGDOUBLE (not LONG_DOUBLE) to
@@ -271,16 +275,29 @@ DEF_FUNCTION_TYPE_1 (BT_FN_VOID_VPTR, BT_VOID, BT_VOLATILE_PTR)
 DEF_FUNCTION_TYPE_1 (BT_FN_VOID_PTRPTR, BT_VOID, BT_PTR_PTR)
 DEF_FUNCTION_TYPE_1 (BT_FN_VOID_CONST_PTR, BT_VOID, BT_CONST_PTR)
 DEF_FUNCTION_TYPE_1 (BT_FN_UINT_UINT, BT_UINT, BT_UINT)
+DEF_FUNCTION_TYPE_1 (BT_FN_UINT_INT, BT_UINT, BT_INT)
+DEF_FUNCTION_TYPE_1 (BT_FN_UINT_ULONG, BT_UINT, BT_ULONG)
+DEF_FUNCTION_TYPE_1 (BT_FN_UINT_LONG, BT_UINT, BT_LONG)
+DEF_FUNCTION_TYPE_1 (BT_FN_UINT_PTR, BT_UINT, BT_PTR)
+DEF_FUNCTION_TYPE_1 (BT_FN_ULONG_PTR, BT_ULONG, BT_PTR)
 DEF_FUNCTION_TYPE_1 (BT_FN_ULONG_ULONG, BT_ULONG, BT_ULONG)
 DEF_FUNCTION_TYPE_1 (BT_FN_ULONGLONG_ULONGLONG, BT_ULONGLONG, BT_ULONGLONG)
+DEF_FUNCTION_TYPE_1 (BT_FN_INT8_FLOAT, BT_INT8, BT_FLOAT)
+DEF_FUNCTION_TYPE_1 (BT_FN_INT16_FLOAT, BT_INT16, BT_FLOAT)
+DEF_FUNCTION_TYPE_1 (BT_FN_UINT32_FLOAT, BT_UINT32, BT_FLOAT)
+DEF_FUNCTION_TYPE_1 (BT_FN_UINT16_FLOAT, BT_UINT16, BT_FLOAT)
+DEF_FUNCTION_TYPE_1 (BT_FN_UINT8_FLOAT, BT_UINT8, BT_FLOAT)
 DEF_FUNCTION_TYPE_1 (BT_FN_UINT16_UINT16, BT_UINT16, BT_UINT16)
 DEF_FUNCTION_TYPE_1 (BT_FN_UINT32_UINT32, BT_UINT32, BT_UINT32)
 DEF_FUNCTION_TYPE_1 (BT_FN_UINT64_UINT64, BT_UINT64, BT_UINT64)
+DEF_FUNCTION_TYPE_1 (BT_FN_UINT64_FLOAT, BT_UINT64, BT_FLOAT)
 DEF_FUNCTION_TYPE_1 (BT_FN_BOOL_INT, BT_BOOL, BT_INT)
 DEF_FUNCTION_TYPE_1 (BT_FN_PTR_CONST_PTR, BT_PTR, BT_CONST_PTR)
 DEF_FUNCTION_TYPE_1 (BT_FN_CONST_PTR_CONST_PTR, BT_CONST_PTR, BT_CONST_PTR)
 DEF_FUNCTION_TYPE_1 (BT_FN_BND_CONST_PTR, BT_BND, BT_CONST_PTR)
 DEF_FUNCTION_TYPE_1 (BT_FN_CONST_PTR_BND, BT_CONST_PTR, BT_BND)
+DEF_FUNCTION_TYPE_1 (BT_FN_UINT16_UINT32, BT_UINT16, BT_UINT32)
+DEF_FUNCTION_TYPE_1 (BT_FN_UINT32_UINT16, BT_UINT32, BT_UINT16)
 
 DEF_POINTER_TYPE (BT_PTR_FN_VOID_PTR, BT_FN_VOID_PTR)
 
@@ -301,18 +318,52 @@ DEF_FUNCTION_TYPE_2 (BT_FN_INT_CONST_STRING_FILEPTR,
 		     BT_INT, BT_CONST_STRING, BT_FILEPTR)
 DEF_FUNCTION_TYPE_2 (BT_FN_INT_INT_FILEPTR,
 		     BT_INT, BT_INT, BT_FILEPTR)
+DEF_FUNCTION_TYPE_2 (BT_FN_UINT16_UINT16_UINT16,
+		     BT_UINT16, BT_UINT16, BT_UINT16)
+DEF_FUNCTION_TYPE_2 (BT_FN_INT_PTR_INT,
+     		     BT_INT, BT_PTR, BT_INT)
+DEF_FUNCTION_TYPE_2 (BT_FN_UINT_PTR_UINT,
+     		     BT_UINT, BT_PTR, BT_UINT)
+DEF_FUNCTION_TYPE_2 (BT_FN_LONG_PTR_LONG,
+		     BT_LONG, BT_PTR, BT_LONG)
+DEF_FUNCTION_TYPE_2 (BT_FN_ULONG_PTR_ULONG,
+		     BT_ULONG, BT_PTR, BT_ULONG)
 DEF_FUNCTION_TYPE_2 (BT_FN_VOID_PTRMODE_PTR,
 		     BT_VOID, BT_PTRMODE, BT_PTR)
 DEF_FUNCTION_TYPE_2 (BT_FN_VOID_PTR_PTRMODE,
 		     BT_VOID, BT_PTR, BT_PTRMODE)
+DEF_FUNCTION_TYPE_2 (BT_FN_VOID_UINT64_UINT64,
+     		     BT_VOID, BT_UINT64, BT_UINT64)
 DEF_FUNCTION_TYPE_2 (BT_FN_VOID_VALIST_REF_VALIST_ARG,
 		     BT_VOID, BT_VALIST_REF, BT_VALIST_ARG)
 DEF_FUNCTION_TYPE_2 (BT_FN_LONG_LONG_LONG,
 		     BT_LONG, BT_LONG, BT_LONG)
+DEF_FUNCTION_TYPE_2 (BT_FN_UINT8_UINT8_UINT8,
+		     BT_UINT8, BT_UINT8, BT_UINT8)
+DEF_FUNCTION_TYPE_2 (BT_FN_INT8_INT8_INT8,
+		     BT_INT8, BT_INT8, BT_INT8)
+DEF_FUNCTION_TYPE_2 (BT_FN_INT16_INT16_INT16,
+		     BT_INT16, BT_INT16, BT_INT16)
+DEF_FUNCTION_TYPE_2 (BT_FN_INT_INT_INT,
+		     BT_INT, BT_INT, BT_INT)
+DEF_FUNCTION_TYPE_2 (BT_FN_UINT_FLOAT_UINT,
+		     BT_UINT, BT_FLOAT, BT_UINT)
+DEF_FUNCTION_TYPE_2 (BT_FN_FLOAT_UINT_UINT,
+		     BT_FLOAT, BT_UINT, BT_UINT)
+DEF_FUNCTION_TYPE_2 (BT_FN_ULONG_UINT_UINT,
+		     BT_ULONG, BT_UINT, BT_UINT)
+DEF_FUNCTION_TYPE_2 (BT_FN_ULONG_UINT_PTR,
+		     BT_ULONG, BT_UINT, BT_PTR)
+DEF_FUNCTION_TYPE_2 (BT_FN_ULONG_ULONG_ULONG,
+		     BT_ULONG, BT_ULONG, BT_ULONG)
+DEF_FUNCTION_TYPE_2 (BT_FN_UINT_UINT_UINT,
+		     BT_UINT, BT_UINT, BT_UINT)
 DEF_FUNCTION_TYPE_2 (BT_FN_INT_PTR_CONST_STRING,
 		     BT_INT, BT_PTR, BT_CONST_STRING)
 DEF_FUNCTION_TYPE_2 (BT_FN_VOID_PTR_SIZE,
 		     BT_VOID, BT_PTR, BT_SIZE)
+DEF_FUNCTION_TYPE_2 (BT_FN_VOID_UINT_PTR,
+		     BT_VOID, BT_UINT, BT_PTR)
 DEF_FUNCTION_TYPE_2 (BT_FN_FLOAT_FLOAT_FLOAT,
 		     BT_FLOAT, BT_FLOAT, BT_FLOAT)
 DEF_FUNCTION_TYPE_2 (BT_FN_DOUBLE_DOUBLE_DOUBLE,
@@ -408,6 +459,7 @@ DEF_FUNCTION_TYPE_2 (BT_FN_BOOL_SIZE_CONST_VPTR, BT_BOOL, BT_SIZE,
 		     BT_CONST_VOLATILE_PTR)
 DEF_FUNCTION_TYPE_2 (BT_FN_BOOL_INT_BOOL, BT_BOOL, BT_INT, BT_BOOL)
 DEF_FUNCTION_TYPE_2 (BT_FN_VOID_UINT_UINT, BT_VOID, BT_UINT, BT_UINT)
+DEF_FUNCTION_TYPE_2 (BT_FN_UINT_UINT_PTR, BT_UINT, BT_UINT, BT_PTR)
 DEF_FUNCTION_TYPE_2 (BT_FN_PTR_CONST_PTR_SIZE, BT_PTR, BT_CONST_PTR, BT_SIZE)
 DEF_FUNCTION_TYPE_2 (BT_FN_PTR_CONST_PTR_CONST_PTR, BT_PTR, BT_CONST_PTR, BT_CONST_PTR)
 DEF_FUNCTION_TYPE_2 (BT_FN_VOID_PTRPTR_CONST_PTR, BT_VOID, BT_PTR_PTR, BT_CONST_PTR)
@@ -415,6 +467,8 @@ DEF_FUNCTION_TYPE_2 (BT_FN_VOID_CONST_PTR_SIZE, BT_VOID, BT_CONST_PTR, BT_SIZE)
 DEF_FUNCTION_TYPE_2 (BT_FN_VOID_PTR_BND, BT_VOID, BT_PTR, BT_BND)
 DEF_FUNCTION_TYPE_2 (BT_FN_CONST_PTR_CONST_PTR_CONST_PTR, BT_CONST_PTR, BT_CONST_PTR, BT_CONST_PTR)
 DEF_FUNCTION_TYPE_2 (BT_FN_BND_CONST_PTR_SIZE, BT_BND, BT_CONST_PTR, BT_SIZE)
+DEF_FUNCTION_TYPE_2 (BT_FN_UINT32_UINT64_PTR,
+		     BT_UINT32, BT_UINT64, BT_PTR)
 
 DEF_POINTER_TYPE (BT_PTR_FN_VOID_PTR_PTR, BT_FN_VOID_PTR_PTR)
 
@@ -444,6 +498,20 @@ DEF_FUNCTION_TYPE_3 (BT_FN_INT_FILEPTR_CONST_STRING_VALIST_ARG,
 		     BT_INT, BT_FILEPTR, BT_CONST_STRING, BT_VALIST_ARG)
 DEF_FUNCTION_TYPE_3 (BT_FN_INT_PTR_PTR_PTR,
 		     BT_INT, BT_PTR, BT_PTR, BT_PTR)
+DEF_FUNCTION_TYPE_3 (BT_FN_INT_INT_UINT_UINT,
+		     BT_INT, BT_INT, BT_UINT, BT_UINT)
+DEF_FUNCTION_TYPE_3 (BT_FN_UINT_UINT_UINT_UINT,
+		     BT_UINT, BT_UINT, BT_UINT, BT_UINT)
+DEF_FUNCTION_TYPE_3 (BT_FN_UINT_UINT_UINT_PTR,
+		     BT_UINT, BT_UINT, BT_UINT, BT_PTR)
+DEF_FUNCTION_TYPE_3 (BT_FN_UINT_ULONG_ULONG_UINT,
+		     BT_UINT, BT_ULONG, BT_ULONG, BT_UINT)
+DEF_FUNCTION_TYPE_3 (BT_FN_ULONG_ULONG_ULONG_ULONG,
+		     BT_ULONG, BT_ULONG, BT_ULONG, BT_ULONG)
+DEF_FUNCTION_TYPE_3 (BT_FN_LONG_LONG_UINT_UINT,
+		     BT_LONG, BT_LONG, BT_UINT, BT_UINT)
+DEF_FUNCTION_TYPE_3 (BT_FN_ULONG_ULONG_UINT_UINT,
+		     BT_ULONG, BT_ULONG, BT_UINT, BT_UINT)
 DEF_FUNCTION_TYPE_3 (BT_FN_STRING_CONST_STRING_CONST_STRING_INT,
 		     BT_STRING, BT_CONST_STRING, BT_CONST_STRING, BT_INT)
 DEF_FUNCTION_TYPE_3 (BT_FN_FLOAT_FLOAT_FLOAT_FLOAT,
@@ -512,6 +580,10 @@ DEF_FUNCTION_TYPE_3 (BT_FN_BOOL_ULONG_ULONG_ULONGPTR, BT_BOOL, BT_ULONG,
 		     BT_ULONG, BT_PTR_ULONG)
 DEF_FUNCTION_TYPE_3 (BT_FN_BOOL_ULONGLONG_ULONGLONG_ULONGLONGPTR, BT_BOOL,
 		     BT_ULONGLONG, BT_ULONGLONG, BT_PTR_ULONGLONG)
+DEF_FUNCTION_TYPE_3 (BT_FN_VOID_UINT32_UINT64_PTR,
+		     BT_VOID, BT_UINT32, BT_UINT64, BT_PTR)
+DEF_FUNCTION_TYPE_3 (BT_FN_VOID_UINT32_UINT32_PTR,
+		     BT_VOID, BT_UINT32, BT_UINT32, BT_PTR)
 
 DEF_FUNCTION_TYPE_4 (BT_FN_SIZE_CONST_PTR_SIZE_SIZE_FILEPTR,
 		     BT_SIZE, BT_CONST_PTR, BT_SIZE, BT_SIZE, BT_FILEPTR)
@@ -523,6 +595,12 @@ DEF_FUNCTION_TYPE_4 (BT_FN_PTR_PTR_CONST_PTR_SIZE_SIZE,
 		     BT_PTR, BT_PTR, BT_CONST_PTR, BT_SIZE, BT_SIZE)
 DEF_FUNCTION_TYPE_4 (BT_FN_PTR_PTR_INT_SIZE_SIZE,
 		     BT_PTR, BT_PTR, BT_INT, BT_SIZE, BT_SIZE)
+DEF_FUNCTION_TYPE_4 (BT_FN_UINT_UINT_UINT_UINT_UINT,
+		     BT_UINT, BT_UINT, BT_UINT, BT_UINT, BT_UINT)
+DEF_FUNCTION_TYPE_4 (BT_FN_UINT_FLOAT_FLOAT_FLOAT_FLOAT,
+		     BT_UINT, BT_FLOAT, BT_FLOAT, BT_FLOAT, BT_FLOAT)
+DEF_FUNCTION_TYPE_4 (BT_FN_ULONG_ULONG_ULONG_UINT_UINT,
+		     BT_ULONG, BT_ULONG, BT_ULONG, BT_UINT, BT_UINT)
 DEF_FUNCTION_TYPE_4 (BT_FN_STRING_STRING_CONST_STRING_SIZE_SIZE,
 		     BT_STRING, BT_STRING, BT_CONST_STRING, BT_SIZE, BT_SIZE)
 DEF_FUNCTION_TYPE_4 (BT_FN_INT_FILEPTR_INT_CONST_STRING_VALIST_ARG,

gcc/builtins.def

@@ -1000,5 +1000,48 @@ DEF_GCC_BUILTIN (BUILT_IN_LINE, "LINE", BT_FN_INT, ATTR_NOTHROW_LEAF_LIST)
 /* Pointer Bounds Checker builtins.  */
 #include "chkp-builtins.def"
 
+/* Do not expose the BRIG builtins by default gcc-wide, but only privately in
+   the BRIG FE as long as there are no references for them in the middle end
+   or any of the upstream backends.  */
+
+#ifndef DEF_HSAIL_BUILTIN
+#define DEF_HSAIL_BUILTIN(ENUM, HSAIL_OPCODE, HSAIL_TYPE, NAME, TYPE, ATTRS) \
+  DEF_BUILTIN_STUB (ENUM, "__builtin_" NAME)
+#endif
+
+/* HSAIL atomic builtins do not have separate identifying opcodes.  */
+
+#ifndef DEF_HSAIL_ATOMIC_BUILTIN
+#define DEF_HSAIL_ATOMIC_BUILTIN(ENUM, ATOMIC_OPCODE, HSAIL_TYPE, NAME, \
+	TYPE, ATTRS) \
+  DEF_BUILTIN_STUB (ENUM, "__builtin_" NAME)
+#endif
+
+/* HSAIL saturating arithmetics builtins.  */
+
+#ifndef DEF_HSAIL_SAT_BUILTIN
+#define DEF_HSAIL_SAT_BUILTIN(ENUM, BRIG_OPCODE, HSAIL_TYPE, NAME, \
+	TYPE, ATTRS) \
+  DEF_BUILTIN_STUB (ENUM, "__builtin_" NAME)
+#endif
+
+/* HSAIL builtins used internally by the frontend.  */
+
+#ifndef DEF_HSAIL_INTR_BUILTIN
+#define DEF_HSAIL_INTR_BUILTIN(ENUM, NAME, TYPE, ATTRS) \
+  DEF_BUILTIN_STUB (ENUM, "__builtin_" NAME)
+#endif
+
+/* HSAIL saturated conversions.  */
+
+#ifndef DEF_HSAIL_CVT_ZEROI_SAT_BUILTIN
+#define DEF_HSAIL_CVT_ZEROI_SAT_BUILTIN(ENUM, HSAIL_DEST_TYPE, HSAIL_SRC_TYPE, \
+  NAME, TYPE, ATTRS) \
+  DEF_BUILTIN_STUB (ENUM, "__builtin_" NAME)
+#endif
+
+/* HSAIL/BRIG frontend builtins.  */
+#include "brig-builtins.def"
+
 #undef DEF_BUILTIN_CHKP
 #undef DEF_BUILTIN

gcc/doc/frontends.texi

@@ -16,7 +16,7 @@
 GCC stands for ``GNU Compiler Collection''.  GCC is an integrated
 distribution of compilers for several major programming languages.  These
 languages currently include C, C++, Objective-C, Objective-C++,
-Fortran, Ada, and Go.
+Fortran, Ada, Go, and BRIG (HSAIL).
 
 The abbreviation @dfn{GCC} has multiple meanings in common use.  The
 current official meaning is ``GNU Compiler Collection'', which refers

gcc/doc/install.texi

@@ -2665,7 +2665,10 @@ separately.
 
 Second, you must have the testing tools installed.  This includes
 @uref{http://www.gnu.org/software/dejagnu/,,DejaGnu}, Tcl, and Expect;
-the DejaGnu site has links to these.
+the DejaGnu site has links to these. For running the BRIG frontend
+tests, a tool to assemble the binary BRIGs from HSAIL text,
+@uref{https://github.com/HSAFoundation/HSAIL-Tools/,,HSAILasm} must
+be installed.
 
 If the directories where @command{runtest} and @command{expect} were
 installed are not in the @env{PATH}, you may need to set the following

gcc/doc/invoke.texi

@@ -1338,6 +1338,9 @@ traditional preprocessor).
 @item @var{file}.go
 Go source code.
 
+@item @var{file}.brig
+BRIG files (binary representation of HSAIL).
+
 @item @var{file}.ads
 Ada source code file that contains a library unit declaration (a
 declaration of a package, subprogram, or generic, or a generic
@@ -1386,6 +1389,7 @@ assembler  assembler-with-cpp
 ada
 f77  f77-cpp-input f95  f95-cpp-input
 go
+brig
 @end smallexample
 
 @item -x none

gcc/doc/standards.texi

@@ -301,6 +301,14 @@ available online, see @uref{http://gcc.gnu.org/readings.html}
 As of the GCC 4.7.1 release, GCC supports the Go 1 language standard,
 described at @uref{http://golang.org/doc/go1.html}.
 
+@section HSA Intermediate Language (HSAIL)
+
+GCC can compile the binary representation (BRIG) of the HSAIL text format as
+described in HSA Programmer's Reference Manual version 1.0.1. This
+capability is typically utilized to implement the HSA runtime API's HSAIL 
+finalization extension for a gcc supported processor. HSA standards are
+freely available at @uref{http://www.hsafoundation.com/standards/}.
+
 @section References for Other Languages
 
 @xref{Top, GNAT Reference Manual, About This Guide, gnat_rm,

gcc/testsuite/ChangeLog

+2017-01-24  Pekka Jääskeläinen <pekka@parmance.com>
+	    Martin Jambor  <mjambor@suse.cz>
+
+	* lib/brig-dg.exp: New file.
+	* lib/brig.exp: Likewise.
+	* brig.dg/README: Likewise.
+	* brig.dg/dg.exp: Likewise.
+	* brig.dg/test/gimple/alloca.hsail: Likewise.
+	* brig.dg/test/gimple/atomics.hsail: Likewise.
+	* brig.dg/test/gimple/branches.hsail: Likewise.
+	* brig.dg/test/gimple/fbarrier.hsail: Likewise.
+	* brig.dg/test/gimple/function_calls.hsail: Likewise.
+	* brig.dg/test/gimple/kernarg.hsail: Likewise.
+	* brig.dg/test/gimple/mem.hsail: Likewise.
+	* brig.dg/test/gimple/mulhi.hsail: Likewise.
+	* brig.dg/test/gimple/packed.hsail: Likewise.
+	* brig.dg/test/gimple/smoke_test.hsail: Likewise.
+	* brig.dg/test/gimple/variables.hsail: Likewise.
+	* brig.dg/test/gimple/vector.hsail: Likewise.
+
 2017-01-24  Jakub Jelinek  <jakub@redhat.com>
 
 	* g++.dg/asan/asan_test.C: Enable on all *-*-linux* targets that

gcc/testsuite/brig.dg/README

+BRIG (HSAIL) frontend test cases
+--------------------------------
+
+The suite consists of "smoke tests" that test several features of
+the compilation and regression tests, but is not an exhaustive test
+suite for all HSAIL instructions. The HSA PRM conformance suite
+is supposed to be used for that.
+
+HSAILasm is required for converting the text HSAIL files to BRIGs
+which the compiler consumes. It can be built from
+https://github.com/HSAFoundation/HSAIL-Tools
+

gcc/testsuite/brig.dg/dg.exp

+#   Copyright (C) 2009-2014 Free Software Foundation, Inc.
+
+# This program 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 of the License, or
+# (at your option) any later version.
+#
+# This program is distributed in the hope that it will be useful,
+# but WITHOUT ANY WARRANTY; without even the implied warranty of
+# MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+# GNU General Public License for more details.
+#
+# You should have received a copy of the GNU General Public License
+# along with GCC; see the file COPYING3.  If not see
+# <http://www.gnu.org/licenses/>.
+
+# GCC testsuite that uses the `dg.exp' driver.
+
+load_lib brig-dg.exp
+
+# Initialize `dg'.
+dg-init
+
+if [expr [llength [auto_execok HSAILasm]] > 0] {
+    dg-runtest [find $srcdir/$subdir *.hsail] "" ""
+} else {
+    unsupported "All BRIG FE tests require HSAILasm in PATH."
+}
+
+# All done.
+dg-finish

gcc/testsuite/brig.dg/test/gimple/alloca.hsail

+module &module:1:0:$full:$large:$default;
+
+/* Tests for alloca. */
+
+/* { dg-do compile } */
+/* { dg-options "-fdump-tree-gimple" } */
+
+prog function &subfunction(arg_u32 %return_value)() {
+     alloca_align(1)_u32 $s2, 256;
+     st_arg_u32 $s2, [%return_value];
+     ret;
+};
+
+prog kernel &kernel(kernarg_u64 %input_ptr, kernarg_u64 %output_ptr)
+{
+        ld_kernarg_u64 $d0, [%input_ptr];
+        ld_global_u32 $s0, [$d0];
+
+	alloca_align(256)_u32 $s1, 16;
+	{
+		arg_u32 %return_value;
+		call &subfunction(%return_value)();
+		ld_arg_u32 $s1, [%return_value];
+	}
+        ld_kernarg_u64 $d1, [%output_ptr];
+        st_global_u32 $s1, [$d0];
+};
+
+/* { dg-final { scan-tree-dump "s2 = __builtin___hsail_alloca \\\(256, 1, __context\\\);" "gimple" } } */
+
+/* { dg-final { scan-tree-dump "s1 = __builtin___hsail_alloca \\\(16, 256, __context\\\);" "gimple" } } */
+
+
+/* Both functions should have an alloca frame push and pop. */
+/* { dg-final { scan-tree-dump-times "__builtin___hsail_alloca_push_frame \\\(__context\\\);" 2 "gimple" } } */
+
+/* { dg-final { scan-tree-dump-times "__builtin___hsail_alloca_pop_frame \\\(__context\\\);" 2 "gimple" } } */

gcc/testsuite/brig.dg/test/gimple/atomics.hsail

+module &module:1:0:$full:$large:$default;
+
+/* Test for atomic instructions. */
+
+/* { dg-do compile } */
+/* { dg-options "-fdump-tree-original" } */
+
+prog kernel &Kernel(kernarg_u64 %input_ptr, kernarg_u64 %output_ptr)
+{
+        ld_kernarg_u64 $d0, [%input_ptr];
+
+	atomic_ld_global_rlx_system_b32 $s0, [$d0];
+	atomic_add_global_rlx_system_u32 $s1, [$d0 + 4], $s0;
+
+        ld_kernarg_u64 $d0, [%output_ptr];
+        atomicnoret_st_global_rlx_system_b32 [$d0], $s2;
+
+	atomicnoret_min_global_rlx_system_u32 [$d0 + 4], $s1;
+
+        ret;
+};
+
+/* The atomic loads are implemented by casting to an atomic pointer. */
+/* { dg-final { scan-tree-dump "s0 = VIEW_CONVERT_EXPR<unsigned int>\\\(\\\*\\\(atomic unsigned int \\\*\\\)" "original"} } */
+
+/* The atomic add should call a gcc builtin. */
+/* { dg-final { scan-tree-dump "= __sync_fetch_and_add_4 \\\(" "original"} } */
+
+/* The atomic stores are implemented by casting to an atomic pointer. */
+/* { dg-final { scan-tree-dump "\\\*\\\(atomic unsigned int \\\*\\\) d0 = s2;" "original"} } */
+
+/* The atomic min is implemented by a custom builtin. */
+/* { dg-final { scan-tree-dump "builtin_out.\[0-9\]+ = __builtin___hsail_atomic_min_u32 \\\(" "original"} } */

gcc/testsuite/brig.dg/test/gimple/branches.hsail

+module &module:1:0:$full:$large:$default;
+
+/* Test different style of branches. */
+
+/* { dg-do compile } */
+/* { dg-options "-fdump-tree-gimple" } */
+
+prog kernel &Kernel(kernarg_u64 %input_ptr, kernarg_u64 %output_ptr)
+{
+        ld_kernarg_u64 $d0, [%input_ptr];
+        ld_global_u64 $d1, [$d0];
+        ld_global_u64 $d2, [$d0 + 8];
+
+        ld_global_u32 $s0, [$d0 + 16];
+        ld_global_u32 $s1, [$d0 + 20];
+
+	sbr_width(all)_u32 $s1 [@case0, @case1, @case2];
+@case0:
+        st_global_u64 0, [$d0];
+	br @out;
+@case1:
+        st_global_u64 1, [$d0];
+	br @out;
+@case2:
+        st_global_u64 2, [$d0];
+@out:
+	cmp_eq_u32_u32 $s2, $s1, $s0;
+	cvt_b1_u32 $c0, $s2;
+
+	cbr_width(all)_b1 $c0, @true_branch;
+@false_branch:
+        st_global_u64 $d1, [$d0];
+
+@true_branch:
+        ld_kernarg_u64 $d0, [%output_ptr];
+
+        st_global_u32 $s2, [$d0 + 8];
+	br @skip;
+        st_global_u32 $s3, [$d0 + 12];
+
+@skip:
+        ret;
+};
+
+/* sbr is converted to a switch */
+/* { dg-final { scan-tree-dump "switch \\\(s1\\\) <default: <D.\[0-9\]+>, case 0: <D.\[0-9\]+>, case 1: <D.\[0-9\]+>, case 2: <D.\[0-9\]+>>" "gimple"} } */
+
+/* br @out converted to gotos */
+/* { dg-final { scan-tree-dump-times "goto @out" 2 "gimple"} } */ 
+
+/* the comparison instruction */
+/* { dg-final { scan-tree-dump "c0 = s2 != 0;" "gimple" } } */
+
+/* cbr to an if clause */
+/* { dg-final { scan-tree-dump "if \\\(c0 != 0\\\) goto @true_branch; else goto <D.\[0-9\]+>;" "gimple" } } */
+
+/* br @skip converted to a goto */
+/* { dg-final { scan-tree-dump "goto @skip" "gimple"} } */ 

gcc/testsuite/brig.dg/test/gimple/fbarrier.hsail

+module &module:1:0:$full:$large:$default;
+
+/* Tests for fbarrier. */
+
+/* { dg-do compile } */
+/* { dg-options "-fdump-tree-gimple" } */
+
+fbarrier &fb_module_scope;
+
+prog function &subfunction(arg_u32 %return_value)() {
+
+     workitemflatabsid_u32 $s3;
+     cvt_b1_u32 $c1, $s3;
+     cbr_width(all)_b1 $c1, @skip_fbar;
+     waitfbar &fb_module_scope;
+@skip_fbar:
+
+     st_arg_u32 $s3, [%return_value];
+     ret;
+};
+
+prog kernel &kernel(kernarg_u64 %input_ptr, kernarg_u64 %output_ptr)
+{
+	fbarrier %fb_func_scope;
+
+        ld_kernarg_u64 $d0, [%input_ptr];
+        ld_global_u32 $s0, [$d0];
+
+	workitemflatabsid_u32 $s1;
+	cvt_b1_u32 $c1, $s1;
+	cbr_width(all)_b1 $c1, @skip_init;
+
+	initfbar &fb_module_scope;
+	initfbar %fb_func_scope;
+
+        joinfbar &fb_module_scope;
+
+@skip_init:
+	barrier_width(all);
+
+        joinfbar %fb_func_scope;
+
+	{
+		arg_u32 %return_value;
+		call &subfunction(%return_value)();
+		ld_arg_u32 $s1, [%return_value];
+	}
+	arrivefbar %fb_func_scope;
+
+        ld_kernarg_u64 $d1, [%output_ptr];
+        st_global_u32 $s1, [$d0];
+
+	workitemflatabsid_u32 $s1;
+	cvt_b1_u32 $c0, $s1;
+	cbr_width(all)_b1 $c0, @skip_fini;
+
+	releasefbar &fb_module_scope;
+	releasefbar %fb_func_scope;
+
+@skip_fini:
+
+};
+/* fbarriers are allocated from the group memory in the order of 
+   appearance. The current implementation allocates 32B per fbarrier. */
+
+/* { dg-final { scan-tree-dump "__hsail_waitfbar \\\(0, __context\\\);" "gimple"} } */
+/* { dg-final { scan-tree-dump "__hsail_initfbar \\\(0, __context\\\);" "gimple"} } */
+/* { dg-final { scan-tree-dump "__hsail_initfbar \\\(32, __context\\\);" "gimple"} } */
+/* { dg-final { scan-tree-dump "__hsail_joinfbar \\\(0, __context\\\);" "gimple"} } */
+/* { dg-final { scan-tree-dump "@skip_init:\[\n ]+__builtin___hsail_barrier \\\(__context\\\);\[\n ]+__builtin___hsail_joinfbar \\\(32, __context\\\);" "gimple"} } */
+
+/* { dg-final { scan-tree-dump "__hsail_arrivefbar \\\(32, __context\\\);" "gimple"} } */
+
+/* { dg-final { scan-tree-dump "__hsail_releasefbar \\\(0, __context\\\);\[\n ]+__builtin___hsail_releasefbar \\\(32, __context\\\);" "gimple"} } */

gcc/testsuite/brig.dg/test/gimple/function_calls.hsail

+module &module:1:0:$full:$large:$default;
+
+/* Function calls and argument passing. */
+
+/* { dg-do compile } */
+/* { dg-options "-fdump-tree-gimple" } */
+
+prog function &subfunction(arg_u32 %return_value)(arg_f32 %float_arg, arg_f64 %double_arg, arg_f16 %half_arg) {
+     ld_arg_f32 $s0, [%float_arg];
+     cvt_u32_f32 $s0, $s0;
+
+     ld_arg_f64 $d0, [%double_arg];
+     cvt_u32_f64 $s1, $d0;
+
+     ld_arg_f16 $s2, [%half_arg];
+     cvt_u32_f16 $s2, $s2;
+
+     add_u32 $s3, $s0, $s1;
+     add_u32 $s3, $s3, $s2;
+
+     st_arg_u32 $s3, [%return_value];
+     ret;
+};
+
+prog kernel &kernel(kernarg_u64 %input_ptr, kernarg_u64 %output_ptr)
+{
+        ld_kernarg_u64 $d0, [%input_ptr];
+        ld_global_u32 $s0, [$d0];
+	{
+		arg_f32 %float_arg;
+		arg_f64 %double_arg;
+		arg_f16 %half_arg;
+		arg_u32 %return_value;
+
+		st_arg_f32 12.0f, [%float_arg];
+		st_arg_f64 640.0d, [%double_arg];
+		st_arg_f16 12.0h, [%half_arg];
+
+		call &subfunction(%return_value)(%float_arg, %double_arg, %half_arg);
+
+		ld_arg_u32 $s1, [%return_value];
+	}	
+        ld_kernarg_u64 $d1, [%output_ptr];
+        st_global_u32 $s1, [$d0];
+};
+
+/* The generated function call should have the incoming arguments and three hidden arguments. */
+
+/* { dg-final { scan-tree-dump "_\[0-9\]+ = subfunction \\\(_kernel.float_arg.\[_0-9\]+, _kernel.double_arg.\[_0-9\]+, _kernel.half_arg.\[_0-9\]+, __context, __group_base_addr, __private_base_addr\\\);" "gimple"} } */
+
+/* The callee should refer directly to the scalar arguments when it reads them. */
+/* { dg-final { scan-tree-dump "= float_arg;" "gimple"} } */
+/* { dg-final { scan-tree-dump "= double_arg;" "gimple"} } */
+/* { dg-final { scan-tree-dump "= half_arg;" "gimple"} } */
+
+/* The return value is stored to a temporary before returned. */
+/* { dg-final { scan-tree-dump "_retvalue_temp = s3;" "gimple"} } */
+/* { dg-final { scan-tree-dump "D.\[0-9\]+ = _retvalue_temp;" "gimple"} } */
+/* { dg-final { scan-tree-dump "return D.\[0-9\]+;" "gimple"} } */

gcc/testsuite/brig.dg/test/gimple/kernarg.hsail

+module &module:1:0:$full:$large:$default;
+
+/* Tests for kernarg addressing modes. */
+
+/* { dg-do compile } */
+/* { dg-options "-fdump-tree-original" } */
+
+prog kernel &Kernel(kernarg_u64 %input[4], kernarg_u64 %output_ptr, kernarg_u64 %i)
+{
+	ld_kernarg_u64 $d0, [%i];
+	ld_kernarg_u64 $d0, [%input][$d0 + 1];
+
+	ld_kernarg_u64 $d1, [%output_ptr];
+	st_global_u64 $d0, [$d1];
+
+	ret;
+};
+
+/* [%i] */
+/* { dg-final { scan-tree-dump " = \\\*\\\(unsigned long \\\*\\\) \\\(__args \\\+ 40\\\);" "original"} } */
+
+/* [%input][$d0 + 1] */
+/* { dg-final { scan-tree-dump "\\\*\\\(unsigned long \\\*\\\) \\\(\\\(VIEW_CONVERT_EXPR<void \\\*>\\\(\\\(unsigned long\\\) __args\\\) \\\+ \\\(unsigned long\\\) d0\\\) \\\+ 1\\\);" "original"} } */
+
+

gcc/testsuite/brig.dg/test/gimple/mem.hsail

+module &module:1:0:$full:$large:$default;
+
+/* Tests for load/store addressing modes. */
+
+/* { dg-do compile } */
+/* { dg-options "-fdump-tree-original" } */
+
+prog kernel &Kernel(kernarg_u64 %input_ptr, kernarg_u64 %input_ptr2, kernarg_u64 %output_ptr)
+{
+	global_u32 %global_array[4];
+
+	ld_kernarg_u64 $d0, [%input_ptr];
+	ld_kernarg_u64 $d2, [%input_ptr2];
+	ld_global_u32 $s0, [$d0];
+	ld_global_u64 $d1, [$d2 + 4];
+
+	ld_global_u32 $s2, [%global_array][$d1 + 4];
+
+	ld_kernarg_u64 $d0, [%output_ptr];
+	st_global_u32 $s0, [$d0];
+	st_global_u32 $s1, [$d0 + 4];
+	st_global_u32 $s2, [$d0 + 8];
+
+	ret;
+};
+
+/* %input_ptr, %input_ptr2 and %output_ptr accesses should generate offsets to the __args array */
+/* { dg-final { scan-tree-dump "__args;\[\n \]+d0 =" "original"} } */
+/* { dg-final { scan-tree-dump "\\\(__args \\\+ 8\\\);\[\n \]+d2 =" "original"} } */
+/* { dg-final { scan-tree-dump "\\\(__args \\\+ 16\\\);\[\n \]+d0 =" "original"} } */
+
+/* ld_global_u32 $s0, [$d0] */
+/* { dg-final { scan-tree-dump "\\\*\\\(unsigned int \\\*\\\) d0;\[\n \]+s0 =" "original"} } */
+
+/* ld_global_u64 $d1, [$d2 + 4] pointer arithmetics*/
+/* { dg-final { scan-tree-dump "d2 \\\+ 4\\\);\[\n \]+d1 = " "original"} } */
+
+/* ld_global_u32 $s2, [%global_array][$d1 + 4]; is the most complex form */
+/* { dg-final { scan-tree-dump "\\\(unsigned long\\\) &_Kernel.global_array\\\) \\\+ \\\(unsigned long\\\) d1\\\) \\\+ 4" "original" } } */

gcc/testsuite/brig.dg/test/gimple/mulhi.hsail

+module &module:1:0:$full:$large:$default;
+
+/* Test high part multiplies. */
+
+/* { dg-do compile } */
+/* { dg-options "-fdump-tree-gimple" } */
+
+prog kernel &Kernel(kernarg_u64 %input_ptr, kernarg_u64 %output_ptr)
+{
+        ld_kernarg_u64 $d0, [%input_ptr];
+        ld_global_u64 $d1, [$d0];
+        ld_global_u64 $d2, [$d0 + 8];
+
+        ld_global_u32 $s0, [$d0 + 16];
+        ld_global_u32 $s1, [$d0 + 20];
+
+	mulhi_s32 $s2, $s0, $s1;
+	mulhi_s64 $d2, $d1, $d2;
+
+	mad24hi_s32 $s3, $s0, $s1, $s2;
+	mul24hi_s32 $s3, $s3, $s1;
+
+        ld_kernarg_u64 $d0, [%output_ptr];
+        st_global_u64 $d1, [$d0];
+        st_global_u32 $s2, [$d0 + 8];
+        st_global_u32 $s3, [$d0 + 12];
+
+        ret;
+};
+
+/* All of the hipart mults areImplemented using MULT_HIGHPART_EXPR (h*). */
+/* { dg-final { scan-tree-dump-times " h\\\* " 4 "gimple"} } */
+

gcc/testsuite/brig.dg/test/gimple/packed.hsail

+module &module:1:0:$full:$large:$default;
+
+/* Test for different cases of packed instruction controls. */
+
+/* { dg-do compile } */
+/* { dg-options "-fdump-tree-gimple -fdump-tree-original" } */
+
+prog kernel &Kernel(kernarg_u64 %input_ptr, kernarg_u64 %output_ptr)
+{
+        ld_kernarg_u64 $d0, [%input_ptr];
+        ld_global_b128 $q0, [$d0];
+
+	add_pp_u8x16 $q1, $q0, u8x16(0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15);
+
+	/* Broadcast the 15 as it's the lowest element (pos 0) in the resulting vector. */
+	add_ps_u8x16 $q2, $q1, u8x16(0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15);
+
+	/* Broadcast the lowest element of q1. */
+	add_sp_u8x16 $q3, $q1, $q2;
+
+	/* Perform a scalar computation with the lowest element of both inputs and store it to the lowest element of dest. */
+	add_ss_u8x16 $q4, $q2, $q3;
+
+	/* Saturating arithmetics variations. */
+	add_pp_sat_u8x16 $q5, $q4, u8x16(0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15);
+
+	/* Broadcast the 15 as it's the lowest element (pos 0) in the resulting vector. */
+	add_ps_sat_u8x16 $q6, $q5, u8x16(0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15);
+
+	/* Broadcast the lowest element of q1. */
+	add_sp_sat_u8x16 $q7, $q6, $q5;
+
+	/* Perform a scalar computation with the lowest element of both inputs and store it to the lowest element of dest. */
+	add_ss_sat_u8x16 $q8, $q7, $q6;
+
+	/* Single operand vector computation. */
+	neg_p_s16x8 $q9, $q8;
+
+        ld_kernarg_u64 $d0, [%output_ptr];
+        st_global_b128 $q8, [$d0];
+
+        ret;
+};
+
+/* The b128 load is done using uint128_t*.
+/* { dg-final { scan-tree-dump "q0 = VIEW_CONVERT_EXPR<uint128_t>\\\(mem_read.\[0-9\]+\\\);" "original"} } */
+
+/* Before arithmetics, the uint128_t is casted to a vector datatype. */
+/* { dg-final { scan-tree-dump "<vector\\\(16\\\) unsigned char>\\\(q0\\\) \\\+ \\\{" "original"} } */
+
+/* The u8x16 constant is generated to an array with elements in reverse order */
+/* in comparison to the HSAIL syntax. */
+/* { dg-final { scan-tree-dump "\\\+ { 15, 14, 13, 12, 11, 10, 9, 8, 7, 6, 5, 4, 3, 2, 1, 0 }" "original"} } */
+
+/* After arithmetics, the vector DT is casted back to a uint128_t. */
+/* { dg-final { scan-tree-dump "q1 = VIEW_CONVERT_EXPR<uint128_t>" "original"} } */
+
+/* Broadcasted the constant vector's lowest element and summed it up in the next line. */
+/* { dg-final { scan-tree-dump "= { 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15 };\[\n \]+_\[0-9\]+ = _\[0-9\]+ \\\+ _\[0-9\]+;" "gimple"} } */
+
+/* Broadcasted the registers lowest element via a VEC_PERM_EXPR that has an all-zeros mask. */
+/* { dg-final { scan-tree-dump "VEC_PERM_EXPR <_\[0-9\]+, _\[0-9\]+, { 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0 }>;" "gimple" } } */
+
+/* For the add_ss we assume performing the computation over the whole vector is cheaper than */
+/* extracting the scalar and performing a scalar operation. This aims to stay in the vector
+/* datapath as long as possible. */
+/* { dg-final { scan-tree-dump "_\[0-9\]+ = VIEW_CONVERT_EXPR<vector\\\(16\\\) unsigned char>\\\(q2\\\);\[\n \]+_\[0-9\]+ = VIEW_CONVERT_EXPR<vector\\\(16\\\) unsigned char>\\\(q3\\\);\[\n \]+_\[0-9\]+ = _\[0-9\]+ \\\+ _\[0-9\]+;" "gimple" } } */
+
+/* Insert the lowest element of the result to the lowest element of the result register. */
+/* { dg-final { scan-tree-dump "= VEC_PERM_EXPR <_\[0-9\]+, new_output.\[0-9\]+_\[0-9\]+, { 16, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15 }>;" "gimple" } } */
+
+/* { dg-final { scan-tree-dump "q4 = VIEW_CONVERT_EXPR<uint128_t>\\\(s_output.\[0-9\]+_\[0-9\]+\\\);" "gimple" } } */
+
+/* The saturating arithmetics are (curently) implemented using scalar builtin calls. */
+/* { dg-final { scan-tree-dump-times "= __builtin___hsail_sat_add_u8" 64 "gimple" } } */
+
+/* A single operand vector instr (neg.) */
+/* { dg-final { scan-tree-dump " = VIEW_CONVERT_EXPR<vector\\\(8\\\) signed short>\\\(q8\\\);\[\n \]+_\[0-9\]+ = -_\[0-9\]+;\[\n \]+" "gimple" } } */

gcc/testsuite/brig.dg/test/gimple/smoke_test.hsail

+module &module:1:0:$full:$large:$default;
+
+/* A basic smoke test. */
+
+/* { dg-do compile } */
+/* { dg-options "-fdump-tree-gimple" } */
+
+prog kernel &Kernel(kernarg_u64 %input_ptr, kernarg_u64 %output_ptr)
+{
+	ld_kernarg_u64 $d0, [%input_ptr];
+	ld_global_u32 $s0, [$d0];
+	ld_global_u32 $s1, [$d0 + 4];
+
+	add_u32 $s2, $s0, $s1;
+	add_u32 $s3, $s0, 4294967295;
+
+	ld_kernarg_u64 $d0, [%output_ptr];
+	st_global_u32 $s2, [$d0];
+	st_global_u32 $s3, [$d0 + 4];
+
+	ret;
+};
+
+prog kernel &KernelWithBarrier(kernarg_u64 %input_ptr, kernarg_u64 %output_ptr)
+{
+	ld_kernarg_u64 $d0, [%input_ptr];
+	ld_global_u32 $s0, [$d0];
+	ld_global_u32 $s1, [$d0 + 4];
+
+	add_u32 $s2, $s0, $s1;
+
+	barrier_width(all);
+
+	add_u32 $s3, $s0, 4294967295;
+
+	ld_kernarg_u64 $d0, [%output_ptr];
+	st_global_u32 $s2, [$d0];
+	st_global_u32 $s3, [$d0 + 4];
+
+	ret;
+};
+
+/* The kernel function itself should have a fingerprint as follows */
+/* _Kernel (unsigned char * __args, void * __context, void * __group_base_addr, void * __private_base_addr) */
+/* { dg-final { scan-tree-dump "_Kernel \\\(unsigned char \\\* __args, void \\\* __context, void \\\* __group_base_addr, void \\\* __private_base_addr\\\)" "gimple"} } */
+
+/* ld_kernarg: mem_read.0 = MEM[(unsigned long *)__args]; */
+/* { dg-final { scan-tree-dump "mem_read.\[0-9\] = MEM\\\[\\\(unsigned long \\\*\\\)__args\\\];" "gimple"} } */
+
+/* The latter ld_global_u32 should be visible as a pointer dereference (after pointer arithmetics on a temporary var): */
+/* mem_read.2 = *D.1691; */
+/* { dg-final { scan-tree-dump "mem_read.\[0-9\] = \\\*\[_0-9\]+;" "gimple"} } */
+
+/* add_u32s should generate +operators */
+/* { dg-final { scan-tree-dump "s2 = s0 \\\+ s1;" "gimple"} } */
+/* { dg-final { scan-tree-dump "s3 = s0 \\\+ 4294967295;" "gimple"} } */
+
+/* The latter st_global_u32 should be visible as a pointer dereference (after pointer arithmetics on a temporary var): */
+/* *D.1694 = s3; */
+/* { dg-final { scan-tree-dump "\\\*\[_0-9\]+ = s3;" "gimple"} } */
+
+/* The return inside the kernel should be generated to a goto to the end of the kernel. */
+/*  goto __kernel_exit; */
+/*  __kernel_exit: */
+/* { dg-final { scan-tree-dump "goto __kernel_exit;" "gimple"} } */
+/* { dg-final { scan-tree-dump "__kernel_exit:" "gimple"} } */
+
+/* Expecting a work item loop because there are no barrier calls. */
+/* { dg-final { scan-tree-dump "if \\\(__local_x < __cur_wg_size_x\\\) goto __wi_loop_x; else goto" "gimple"} } */
+/* { dg-final { scan-tree-dump "if \\\(__local_y < __cur_wg_size_y\\\) goto __wi_loop_y; else goto" "gimple"} } */
+/* { dg-final { scan-tree-dump "if \\\(__local_z < __cur_wg_size_z\\\) goto __wi_loop_z; else goto" "gimple"} } */
+
+/* The launcher should call __hsail_launch_wg_function in this case: */
+/* Kernel (void * __context, void * __group_base_addr) */
+/* { dg-final { scan-tree-dump "Kernel \\\(void \\\* __context, void \\\* __group_base_addr\\\)" "gimple"} } */
+/* { dg-final { scan-tree-dump "__hsail_launch_wg_function \\\(_Kernel, __context, __group_base_addr\\\);" "gimple"} }*/
+
+/* The kernel should have the magic metadata section injected to the ELF. */
+/* TODO: this should be disabled in case not outputting to an ELF. */
+/* Currently ELF is assumed by the brig frontend. Do not check for the context */
+/* as it is likely to change. */
+/* { dg-final { scan-tree-dump "\\\.pushsection phsa\\\.desc\\\.Kernel" "gimple"} }*/
+
+/* The kernel with the barrier call should have the barrier builtin call in between the two summations. */
+/* { dg-final { scan-tree-dump "s2 = s0 \\\+ s1;\[\n \]+__builtin___hsail_barrier \\\(__context\\\);\[\n \]+s3 = s0 \\\+ 4294967295;" "gimple"} } */
+
+/* The kernel with the barrier call's launcher function should call the thread-spawning function. */
+/* { dg-final { scan-tree-dump "__hsail_launch_kernel \\\(_KernelWithBarrier, __context, __group_base_addr\\\);" "gimple" } } */
+
+
+

gcc/testsuite/brig.dg/test/gimple/variables.hsail

+module &module:1:0:$full:$large:$default;
+
+/* Tests for different variable scopes and address spaces. */
+
+/* { dg-do compile } */
+/* { dg-options "-fdump-tree-gimple" } */
+
+prog align(256) private_u32 &prog_private;
+private_u32 &mod_private;
+
+prog group_u32 &prog_group;
+group_u32 &mod_group;
+
+prog global_u32 &prog_global;
+global_u32 &mod_global;
+
+decl prog global_u32 &prog_global_host_def;
+
+prog readonly_u32 &prog_readonly;
+readonly_u32 &mod_readonly;
+
+prog function &subfunction(arg_u32 %return_value)(arg_u32 %arg) {
+
+     private_u32 %func_private;
+     group_u32 %func_group;
+     align(256) global_u32 %func_global;
+     readonly_u32 %func_readonly;
+
+     ld_private_u32 $s200, [%func_private];
+     st_private_u32 $s200, [&prog_private];
+
+     ld_group_u32 $s203, [%func_group];
+     st_group_u32 $s203, [&prog_group];
+
+     ld_global_u32 $s204, [%func_global];
+     st_global_u32 $s204, [&prog_global];
+
+     ld_readonly_u32 $s205, [%func_readonly];
+     st_global_u32 $s205, [%func_global];
+
+     st_arg_u32 $s2, [%return_value];
+     ret;
+};
+
+prog kernel &kernel(kernarg_u64 %input_ptr, kernarg_u64 %output_ptr)
+{
+	private_u32 %kern_private;
+	group_u32 %kern_group;
+	global_u32 %kern_global;
+	readonly_u32 %kern_readonly;
+
+        ld_kernarg_u64 $d0, [%input_ptr];
+        ld_global_u32 $s0, [$d0];
+
+	ld_private_u32 $s2, [&prog_private];
+	st_private_u32 $s2, [%kern_private];
+	ld_private_u32 $s3, [&mod_private];
+	st_private_u32 $s3, [&prog_private];
+
+	ld_group_u32 $s4, [&prog_group];
+	st_group_u32 $s4, [%kern_group];
+	ld_group_u32 $s5, [&mod_group];
+	st_group_u32 $s5, [&prog_group];
+
+	ld_global_u32 $s6, [&prog_global];
+	st_global_u32 $s6, [%kern_global];
+	ld_global_u32 $s7, [&mod_global];
+	st_global_u32 $s7, [&prog_global];
+
+	ld_readonly_u32 $s8, [&prog_readonly];
+	st_global_u32 $s8, [%kern_global];
+	ld_readonly_u32 $s9, [&mod_readonly];
+	st_global_u32 $s9, [&prog_global];
+
+	ld_readonly_u32 $s10, [%kern_readonly];
+	st_global_u32 $s10, [%kern_global];
+	ld_readonly_u32 $s11, [%kern_readonly];
+	st_global_u32 $s11, [&prog_global_host_def];
+
+	{
+		arg_u32 %arg;
+		arg_u32 %return_value;
+		st_arg_u32 $s1, [%arg];
+		call &subfunction(%return_value)(%arg);
+		ld_arg_u32 $s1, [%return_value];
+	}
+        ld_kernarg_u64 $d1, [%output_ptr];
+        st_global_u32 $s1, [$d0];
+};
+
+/* Private variable offsets assigned in the order of their appearance */
+/*
+ prog_private @0	(align 256) -> until 254 to ensure all WIs
+ mod_private  @256	               have their chunks aligned
+ func_private @260
+ kern_private @264
+*/
+
+/* Group variable offsets assigned in the order of their appearance */
+/*
+ prog_group @0		(2)
+ mod_group  @4		(4)
+ func_group @8		(1)
+ kern_group @12		(3)
+*/
+
+/* { dg-final { scan-tree-dump "\\\+ 8;.*\\\+ 12;.*\\\+ 4;" "gimple" } } */
+
+/* The "mangling" of the global and readonly vars. */
+/* { dg-final { scan-tree-dump "\[ \]*prog_global = s204;" "gimple" } } */
+
+/* { dg-final { scan-tree-dump "\.module.mod_global;" "gimple" } } */
+
+/* Host defined variables need indirect access as the address is
+   known only at run time. */
+/* { dg-final { scan-tree-dump "\\\*\\\__phsa.host_def.prog_global_host_def.\[0-9\]+_\[0-9\]+ = s11;" "gimple" } } */
+
+/* { dg-final { scan-tree-dump "\.subfunction.func_global;" "gimple" } } */
+/* { dg-final { scan-tree-dump "\.subfunction.func_readonly;" "gimple" } } */
+
+/* { dg-final { scan-tree-dump "kernel.kern_global" "gimple" } } */
+/* { dg-final { scan-tree-dump "kernel.kern_readonly" "gimple" } } */
+
+

gcc/testsuite/brig.dg/test/gimple/vector.hsail

+module &module:1:0:$full:$large:$default;
+
+/* A test for vector operands. */
+
+/* { dg-do compile } */
+/* { dg-options "-fdump-tree-original" } */
+
+prog kernel &Kernel(kernarg_u64 %input_ptr, kernarg_u64 %output_ptr)
+{
+        ld_kernarg_u64 $d0, [%input_ptr];
+	ld_v2_global_f32 ($s0, $s1), [$d0];
+	ld_v3_global_f32 ($s2, $s3, $s4), [$d0 + 8];
+	ld_v4_global_f32 ($s5, $s6, $s7, $s8), [$d0 + 20];
+
+	add_f32 $s9, $s0, $s1;
+	combine_v2_b64_b32 $d2, ($s1, $s0);
+	combine_v2_b64_b32 $d3, ($s2, $s3);
+
+	add_pp_f32x2 $d4, $d2, $d3;
+
+	expand_v2_b32_b64 ($s0, $s3), $d4;
+
+        ld_kernarg_u64 $d1, [%output_ptr];
+        st_v2_global_f32 ($s0, $s1), [$d1];
+        st_v3_global_f32 ($s2, $s3, $s4), [$d1 + 8];
+        st_v4_global_f32 ($s5, $s6, $s7, $s8), [$d1 + 20];
+
+        ret;
+};
+
+/* The v2 load is done via casting to a vector datatype ptr. */
+/* { dg-final { scan-tree-dump " = MEM\\\[\\\(vector\\\(2\\\) <float:32> \\\*\\\)" "original"} } */
+
+/* The v3 load is scalarized (at the moment) due to gcc requiring 2's exponent wide vectors. */
+/* { dg-final { scan-tree-dump "s0 = VIEW_CONVERT_EXPR<unsigned int>\\\(BIT_FIELD_REF <mem_read.\[0-9\]+, 32, 0>\\\);\[\n ]+s1 = VIEW_CONVERT_EXPR<unsigned int>\\\(BIT_FIELD_REF <mem_read.\[0-9\]+, 32, 32>\\\);" "original"} } */
+
+/* The v4 load is done via casting to a vector datatype ptr. */
+/* { dg-final { scan-tree-dump " = MEM\\\[\\\(vector\\\(4\\\) <float:32> \\\*\\\)" "original"} } */
+
+/* The combines are generated to vector constructors. */
+/* { dg-final { scan-tree-dump "{s1, s0}" "original"} } */
+/* { dg-final { scan-tree-dump "{s2, s3}" "original"} } */
+
+/* Expands to BIT_FIELD_REFs. */
+/* { dg-final { scan-tree-dump "s0 = BIT_FIELD_REF <d4, 32, 0>;" "original"} } */
+/* { dg-final { scan-tree-dump "s3 = BIT_FIELD_REF <d4, 32, 32>;" "original"} } */
+
+/* The v1 store is done via casting to a vector datatype ptr and constructing a vector from the inputs. */
+/* { dg-final { scan-tree-dump "MEM\\\[\\\(vector\\\(2\\\) <float:32> \\\*\\\)\\\(<float:32> \\\*\\\) d1\\\] = " "original"} } */
+
+/* The v3 store is scalarized (at the moment) due to gcc requiring 2's exponent wide vectors. */
+/* { dg-final { scan-tree-dump "\\\*\\\(<float:32> \\\*\\\) \\\(\\\(sizetype\\\) d1 \\\+ 8\\\) \\\+ 0 = VIEW_CONVERT_EXPR<<float:32>>\\\(s2\\\);" "original"} } */
+/* { dg-final { scan-tree-dump "\\\*\\\(<float:32> \\\*\\\) \\\(\\\(sizetype\\\) d1 \\\+ 8\\\) \\\+ 4 = VIEW_CONVERT_EXPR<<float:32>>\\\(s3\\\);" "original"} } */
+/* { dg-final { scan-tree-dump "\\\*\\\(<float:32> \\\*\\\) \\\(\\\(sizetype\\\) d1 \\\+ 8\\\) \\\+ 8 = VIEW_CONVERT_EXPR<<float:32>>\\\(s4\\\);" "original"} } */
+
+/* The v4 store is done via casting to a vector datatype and constructing a vector from the inputs. */
+/* { dg-final { scan-tree-dump "MEM\\\[\\\(vector\\\(4\\\) <float:32> \\\*\\\)\\\(<float:32> \\\*\\\) \\\(\\\(sizetype\\\) d1 \\\+ 20\\\)\\\] = {VIEW_CONVERT_EXPR<<float:32>>\\\(s5\\\), VIEW_CONVERT_EXPR<<float:32>>\\\(s6\\\), VIEW_CONVERT_EXPR<<float:32>>\\\(s7\\\), VIEW_CONVERT_EXPR<<float:32>>\\\(s8\\\)};" "original"} } */

gcc/testsuite/lib/brig-dg.exp

+#   Copyright (C) 2009-2014 Free Software Foundation, Inc.
+
+# This program 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 of the License, or
+# (at your option) any later version.
+# 
+# This program is distributed in the hope that it will be useful,
+# but WITHOUT ANY WARRANTY; without even the implied warranty of
+# MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+# GNU General Public License for more details.
+# 
+# You should have received a copy of the GNU General Public License
+# along with GCC; see the file COPYING3.  If not see
+# <http://www.gnu.org/licenses/>.
+
+load_lib gcc-dg.exp
+
+# Define brig callbacks for dg.exp.
+
+proc brig-dg-test { prog do_what extra_tool_flags } {
+    set result \
+	[gcc-dg-test-1 brig_target_compile $prog $do_what $extra_tool_flags]
+    
+    set comp_output [lindex $result 0]
+    set output_file [lindex $result 1]
+
+    return [list $comp_output $output_file]
+}

gcc/testsuite/lib/brig.exp

+# Copyright (C) 2009-2016 Free Software Foundation, Inc.
+
+# This program 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 of the License, or
+# (at your option) any later version.
+#
+# This program is distributed in the hope that it will be useful,
+# but WITHOUT ANY WARRANTY; without even the implied warranty of
+# MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+# GNU General Public License for more details.
+#
+# You should have received a copy of the GNU General Public License
+# along with GCC; see the file COPYING3.  If not see
+# <http://www.gnu.org/licenses/>.
+
+load_lib prune.exp
+load_lib gcc-defs.exp
+load_lib timeout.exp
+load_lib target-libpath.exp
+#
+# brig_target_compile -- compile a HSAIL input to BRIG using HSAILasm and then
+#                        compile the BRIG to target ISA using gcc
+
+proc brig_target_compile { source dest type options } {
+    global tmpdir
+    global testname_with_flags
+    if { [file extension $source] == ".hsail" } {
+	# We cannot assume all inputs are .hsail as the dg machinery
+	# calls this for a some c files to check linker plugin support or
+	# similar.
+	set brig_source ${tmpdir}/[file tail ${source}].brig
+	exec HSAILasm $source -o ${brig_source}
+	set source ${brig_source}
+	# Change the testname the .brig.
+	set testname_with_flags [file tail $source]
+    }
+    return [target_compile $source $dest $type $options]
+}
+

include/ChangeLog

+2017-01-24  Pekka Jääskeläinen <pekka@parmance.com>
+	    Martin Jambor  <mjambor@suse.cz>
+
+	* hsa.h: Moved here from libgomp/plugin/hsa.h.
+
 2017-01-04  Richard Earnshaw  <rearnsha@arm.com>
 	    Jiong Wang  <jiong.wang@arm.com>
 

libgomp/ChangeLog

+2017-01-24  Pekka Jääskeläinen <pekka@parmance.com>
+	    Martin Jambor  <mjambor@suse.cz>
+
+	* plugin/hsa.h: Moved to top level include.
+	* plugin/plugin-hsa.c: Chanfgd include of hsa.h accordingly.
+
 2017-01-21  Jakub Jelinek  <jakub@redhat.com>
 
 	PR other/79046