From: Cesar Philippidis <cesar@codesourcery.com>
Date: Thu, 21 Dec 2017 21:40:34 +0000 (-0800)
Subject: Use functional parameters for data mappings in OpenACC child functions
X-Git-Url: http://git.ipfire.org/cgi-bin/gitweb.cgi?a=commitdiff_plain;h=493b01873ee121ba5dd6b60084ee69620b696619;p=thirdparty%2Fgcc.git

Use functional parameters for data mappings in OpenACC child functions

	* Makefile.def: Make libgomp depend on libffi.
	* configure.ac: Likewise.
	* Makefile.in: Regenerate.
	* configure: Regenerate.

	gcc/fortran/
	* types.def: (BF_FN_VOID_INT_INT_OMPFN_SIZE_PTR_PTR_PTR_VAR):
	Define.

	gcc/
	* builtin-types.def (BF_FN_VOID_INT_INT_OMPFN_SIZE_PTR_PTR_PTR_VAR):
	Define.
	* config/nvptx/nvptx.c (nvptx_expand_cmp_swap): Handle PARM_DECLs.
	* omp-builtins.def (BUILD_IN_GOACC_PARALLEL): Call
	GOACC_parallel_keyed_v2.
	* omp-expand.c (expand_omp_target): Update call to
	BUILT_IN_GOACC_PARALLEL.
	* omp-low.c (struct omp_context): Add parm_map member.
	(lookup_parm): New function.
	(build_receiver_ref): Lookup parm_map decls.
	(install_parm_decl): New function.
	(install_var_field): Install parm_map decl for OpenACC parallel region
	data clauses.
	(delete_omp_context): Clean parm_map.
	(scan_sharing_clauses): Install subarray variable mapping into parm_map.
	(create_omp_child_function): Defer creation of child function for
	OpenACC parallel regions.
	(scan_omp_target): Likewise.
	(append_decl_arg): New function.
	(lower_omp_target): Create an child offloaded function using one
	parameter per data mapping for OpenACC parallel regions.
	* tree-ssa-structalias.c (find_func_aliases_for_builtin_call):
	Ignore OpenACC parallel regions.
	(find_func_clobbers): Likewise.
	(ipa_pta_execute): Likewise.

	libgomp/
	* Makefile.am: Add libffi build dependency.
	* configure.ac: Likewise.
	* Makefile.in: Regenerate.
	* config.h.in: Regenerate.
	* configure: Regenerate.
	* libgomp-plugin.h: Define GOMP_OFFLOAD_openacc_exec_params and
	GOMP_OFFLOAD_openacc_async_exec_params.
	* libgomp.h (acc_dispatch_t): Use them here.
	* libgomp.map (GOACC_parallel_keyed_v2): Declare.
	* libgomp_g.h (GOACC_parallel_keyed_v2): Likewise.
	* oacc-host.c (host_openacc_exec_params): New function.
	(host_openacc_async_exec_params): Likewise.
	* oacc-parallel.c (goacc_call_host_fn): Likewise.
	(GOACC_parallel_keyed_internal): Likewise.
	(GOACC_parallel_keyed): Wrapper for GOACC_parallel_keyed_internal.
	(GOACC_parallel_keyed_v2): Likewise.
	* plugin/plugin-nvptx.c (nvptx_exec): Replace CUDeviceptr dp parameter
	with void **kargs.
	(openacc_exec_internal): New function.
	(GOMP_OFFLOAD_openacc_exec_params): New function.
	(GOMP_OFFLOAD_openacc_exec): Update to call openacc_exec_internal.
	(openacc_async_exec_internal): New function.
	(GOMP_OFFLOAD_openacc_async_exec_params): New function.
	(GOMP_OFFLOAD_openacc_async_exec): Update call to
	openacc_async_exec_internal.
	* target.c (gomp_load_plugin_for_device): Handle
	openacc_exec_params and openacc_async_exec_params.
	* testsuite/Makefile.in: Regenerate.
	* testsuite/libgomp.oacc-c-c++-common/combined-directives-1.c:
	Xfail on offloaded targets.

	* Makefile.def: Bootstrap module libffi. Add libffi dependency
	to all-target-libgomp.
	* Makefile.in: Regenerate.
	* configure.ac: Add libffi to bootstrap_target_libs when libgomp
	is bootstrapped.
	* configure: Regenerate.

	gcc/
	* omp-low.c (install_parm_decl): Don't extract identifiers from
	artifical decls.

	gcc/testsuite/
	* c-c++-common/goacc/large_array.c: New test.

(cherry picked from openacc-gcc-7-branch commit
b4dd21b9a1f9f499c613b55225cad689b7928a7f, commit
9ba1d875dcb9412cccdd49138a3525e7adab3e76, commit
762cf3c7890fab15a69494a6480455cd99621d7d, and commit
6585af7290fd79f6cb834a39c2bbf7e1934808b1)

(cherry picked from openacc-gcc-9-branch commit
721ed7ea405710ef0beab3629567dd39cdd22ca3)
---

diff --git a/Makefile.def b/Makefile.def
index 50bd67e5146e..c5cc3125195a 100644
--- a/Makefile.def
+++ b/Makefile.def
@@ -162,7 +162,7 @@ target_modules = { module= libtermcap; no_check=true;
                    missing=maintainer-clean; };
 target_modules = { module= winsup; };
 target_modules = { module= libgloss; no_check=true; };
-target_modules = { module= libffi; no_install=true; };
+target_modules = { module= libffi; bootstrap=true; no_install=true; };
 target_modules = { module= zlib; };
 target_modules = { module= rda; };
 target_modules = { module= libada; };
@@ -553,6 +553,8 @@ dependencies = { module=configure-target-libgo; on=all-target-libstdc++-v3; };
 dependencies = { module=all-target-libgo; on=all-target-libbacktrace; };
 dependencies = { module=all-target-libgo; on=all-target-libffi; };
 dependencies = { module=all-target-libgo; on=all-target-libatomic; };
+dependencies = { module=all-target-libgomp; on=all-target-libffi; };
+dependencies = { module=configure-target-libgomp; on=configure-target-libffi; };
 dependencies = { module=configure-target-libphobos; on=configure-target-libbacktrace; };
 dependencies = { module=configure-target-libphobos; on=configure-target-zlib; };
 dependencies = { module=all-target-libphobos; on=all-target-libbacktrace; };
@@ -573,6 +575,7 @@ dependencies = { module=install-target-libgfortran; on=install-target-libquadmat
 dependencies = { module=install-target-libgfortran; on=install-target-libgcc; };
 dependencies = { module=install-target-libphobos; on=install-target-libatomic; };
 dependencies = { module=install-target-libsanitizer; on=install-target-libstdc++-v3; };
+dependencies = { module=install-target-libgomp; on=install-target-libffi; };
 dependencies = { module=install-target-libsanitizer; on=install-target-libgcc; };
 dependencies = { module=install-target-libvtv; on=install-target-libstdc++-v3; };
 dependencies = { module=install-target-libvtv; on=install-target-libgcc; };
diff --git a/Makefile.in b/Makefile.in
index dbf4fb7c9e3f..e1e2bd73bbf2 100644
--- a/Makefile.in
+++ b/Makefile.in
@@ -1219,7 +1219,9 @@ all-target: maybe-all-target-libphobos
 all-target: maybe-all-target-libtermcap
 all-target: maybe-all-target-winsup
 all-target: maybe-all-target-libgloss
+@if target-libffi-no-bootstrap
 all-target: maybe-all-target-libffi
+@endif target-libffi-no-bootstrap
 all-target: maybe-all-target-zlib
 all-target: maybe-all-target-rda
 all-target: maybe-all-target-libada
@@ -48974,7 +48976,6 @@ configure-target-libffi: stage_current
 @if target-libffi
 maybe-configure-target-libffi: configure-target-libffi
 configure-target-libffi: 
-	@: $(MAKE); $(unstage)
 	@r=`${PWD_COMMAND}`; export r; \
 	s=`cd $(srcdir); ${PWD_COMMAND}`; export s; \
 	echo "Checking multilib configuration for libffi..."; \
@@ -49012,6 +49013,412 @@ configure-target-libffi:
 
 
 
+.PHONY: configure-stage1-target-libffi maybe-configure-stage1-target-libffi
+maybe-configure-stage1-target-libffi:
+@if target-libffi-bootstrap
+maybe-configure-stage1-target-libffi: configure-stage1-target-libffi
+configure-stage1-target-libffi:
+	@[ $(current_stage) = stage1 ] || $(MAKE) stage1-start
+	@$(SHELL) $(srcdir)/mkinstalldirs $(TARGET_SUBDIR)/libffi
+	@r=`${PWD_COMMAND}`; export r; \
+	s=`cd $(srcdir); ${PWD_COMMAND}`; export s; \
+	TFLAGS="$(STAGE1_TFLAGS)"; \
+	echo "Checking multilib configuration for libffi..."; \
+	$(CC_FOR_TARGET) --print-multi-lib > $(TARGET_SUBDIR)/libffi/multilib.tmp 2> /dev/null; \
+	if test -r $(TARGET_SUBDIR)/libffi/multilib.out; then \
+	  if cmp -s $(TARGET_SUBDIR)/libffi/multilib.tmp $(TARGET_SUBDIR)/libffi/multilib.out; then \
+	    rm -f $(TARGET_SUBDIR)/libffi/multilib.tmp; \
+	  else \
+	    rm -f $(TARGET_SUBDIR)/libffi/Makefile; \
+	    mv $(TARGET_SUBDIR)/libffi/multilib.tmp $(TARGET_SUBDIR)/libffi/multilib.out; \
+	  fi; \
+	else \
+	  mv $(TARGET_SUBDIR)/libffi/multilib.tmp $(TARGET_SUBDIR)/libffi/multilib.out; \
+	fi; \
+	test ! -f $(TARGET_SUBDIR)/libffi/Makefile || exit 0; \
+	$(NORMAL_TARGET_EXPORTS) \
+	CFLAGS="$(CFLAGS_FOR_TARGET)"; export CFLAGS; \
+	CXXFLAGS="$(CXXFLAGS_FOR_TARGET)"; export CXXFLAGS; \
+	LIBCFLAGS="$(LIBCFLAGS_FOR_TARGET)"; export LIBCFLAGS;  \
+	echo Configuring stage 1 in $(TARGET_SUBDIR)/libffi; \
+	$(SHELL) $(srcdir)/mkinstalldirs $(TARGET_SUBDIR)/libffi; \
+	cd $(TARGET_SUBDIR)/libffi || exit 1; \
+	case $(srcdir) in \
+	  /* | [A-Za-z]:[\\/]*) topdir=$(srcdir) ;; \
+	  *) topdir=`echo $(TARGET_SUBDIR)/libffi/ | \
+		sed -e 's,\./,,g' -e 's,[^/]*/,../,g' `$(srcdir) ;; \
+	esac; \
+	module_srcdir=libffi; \
+	$(SHELL) $$s/$$module_srcdir/configure \
+	  --srcdir=$${topdir}/$$module_srcdir \
+	  $(TARGET_CONFIGARGS) --build=${build_alias} --host=${target_alias} \
+	  --target=${target_alias} \
+	   \
+	  $(STAGE1_CONFIGURE_FLAGS)
+@endif target-libffi-bootstrap
+
+.PHONY: configure-stage2-target-libffi maybe-configure-stage2-target-libffi
+maybe-configure-stage2-target-libffi:
+@if target-libffi-bootstrap
+maybe-configure-stage2-target-libffi: configure-stage2-target-libffi
+configure-stage2-target-libffi:
+	@[ $(current_stage) = stage2 ] || $(MAKE) stage2-start
+	@$(SHELL) $(srcdir)/mkinstalldirs $(TARGET_SUBDIR)/libffi
+	@r=`${PWD_COMMAND}`; export r; \
+	s=`cd $(srcdir); ${PWD_COMMAND}`; export s; \
+	TFLAGS="$(STAGE2_TFLAGS)"; \
+	echo "Checking multilib configuration for libffi..."; \
+	$(CC_FOR_TARGET) --print-multi-lib > $(TARGET_SUBDIR)/libffi/multilib.tmp 2> /dev/null; \
+	if test -r $(TARGET_SUBDIR)/libffi/multilib.out; then \
+	  if cmp -s $(TARGET_SUBDIR)/libffi/multilib.tmp $(TARGET_SUBDIR)/libffi/multilib.out; then \
+	    rm -f $(TARGET_SUBDIR)/libffi/multilib.tmp; \
+	  else \
+	    rm -f $(TARGET_SUBDIR)/libffi/Makefile; \
+	    mv $(TARGET_SUBDIR)/libffi/multilib.tmp $(TARGET_SUBDIR)/libffi/multilib.out; \
+	  fi; \
+	else \
+	  mv $(TARGET_SUBDIR)/libffi/multilib.tmp $(TARGET_SUBDIR)/libffi/multilib.out; \
+	fi; \
+	test ! -f $(TARGET_SUBDIR)/libffi/Makefile || exit 0; \
+	$(NORMAL_TARGET_EXPORTS) \
+	 \
+	CFLAGS="$(CFLAGS_FOR_TARGET)"; export CFLAGS; \
+	CXXFLAGS="$(CXXFLAGS_FOR_TARGET)"; export CXXFLAGS; \
+	LIBCFLAGS="$(LIBCFLAGS_FOR_TARGET)"; export LIBCFLAGS;  \
+	echo Configuring stage 2 in $(TARGET_SUBDIR)/libffi; \
+	$(SHELL) $(srcdir)/mkinstalldirs $(TARGET_SUBDIR)/libffi; \
+	cd $(TARGET_SUBDIR)/libffi || exit 1; \
+	case $(srcdir) in \
+	  /* | [A-Za-z]:[\\/]*) topdir=$(srcdir) ;; \
+	  *) topdir=`echo $(TARGET_SUBDIR)/libffi/ | \
+		sed -e 's,\./,,g' -e 's,[^/]*/,../,g' `$(srcdir) ;; \
+	esac; \
+	module_srcdir=libffi; \
+	$(SHELL) $$s/$$module_srcdir/configure \
+	  --srcdir=$${topdir}/$$module_srcdir \
+	  $(TARGET_CONFIGARGS) --build=${build_alias} --host=${target_alias} \
+	  --target=${target_alias} \
+	  --with-build-libsubdir=$(HOST_SUBDIR) \
+	  $(STAGE2_CONFIGURE_FLAGS)
+@endif target-libffi-bootstrap
+
+.PHONY: configure-stage3-target-libffi maybe-configure-stage3-target-libffi
+maybe-configure-stage3-target-libffi:
+@if target-libffi-bootstrap
+maybe-configure-stage3-target-libffi: configure-stage3-target-libffi
+configure-stage3-target-libffi:
+	@[ $(current_stage) = stage3 ] || $(MAKE) stage3-start
+	@$(SHELL) $(srcdir)/mkinstalldirs $(TARGET_SUBDIR)/libffi
+	@r=`${PWD_COMMAND}`; export r; \
+	s=`cd $(srcdir); ${PWD_COMMAND}`; export s; \
+	TFLAGS="$(STAGE3_TFLAGS)"; \
+	echo "Checking multilib configuration for libffi..."; \
+	$(CC_FOR_TARGET) --print-multi-lib > $(TARGET_SUBDIR)/libffi/multilib.tmp 2> /dev/null; \
+	if test -r $(TARGET_SUBDIR)/libffi/multilib.out; then \
+	  if cmp -s $(TARGET_SUBDIR)/libffi/multilib.tmp $(TARGET_SUBDIR)/libffi/multilib.out; then \
+	    rm -f $(TARGET_SUBDIR)/libffi/multilib.tmp; \
+	  else \
+	    rm -f $(TARGET_SUBDIR)/libffi/Makefile; \
+	    mv $(TARGET_SUBDIR)/libffi/multilib.tmp $(TARGET_SUBDIR)/libffi/multilib.out; \
+	  fi; \
+	else \
+	  mv $(TARGET_SUBDIR)/libffi/multilib.tmp $(TARGET_SUBDIR)/libffi/multilib.out; \
+	fi; \
+	test ! -f $(TARGET_SUBDIR)/libffi/Makefile || exit 0; \
+	$(NORMAL_TARGET_EXPORTS) \
+	 \
+	CFLAGS="$(CFLAGS_FOR_TARGET)"; export CFLAGS; \
+	CXXFLAGS="$(CXXFLAGS_FOR_TARGET)"; export CXXFLAGS; \
+	LIBCFLAGS="$(LIBCFLAGS_FOR_TARGET)"; export LIBCFLAGS;  \
+	echo Configuring stage 3 in $(TARGET_SUBDIR)/libffi; \
+	$(SHELL) $(srcdir)/mkinstalldirs $(TARGET_SUBDIR)/libffi; \
+	cd $(TARGET_SUBDIR)/libffi || exit 1; \
+	case $(srcdir) in \
+	  /* | [A-Za-z]:[\\/]*) topdir=$(srcdir) ;; \
+	  *) topdir=`echo $(TARGET_SUBDIR)/libffi/ | \
+		sed -e 's,\./,,g' -e 's,[^/]*/,../,g' `$(srcdir) ;; \
+	esac; \
+	module_srcdir=libffi; \
+	$(SHELL) $$s/$$module_srcdir/configure \
+	  --srcdir=$${topdir}/$$module_srcdir \
+	  $(TARGET_CONFIGARGS) --build=${build_alias} --host=${target_alias} \
+	  --target=${target_alias} \
+	  --with-build-libsubdir=$(HOST_SUBDIR) \
+	  $(STAGE3_CONFIGURE_FLAGS)
+@endif target-libffi-bootstrap
+
+.PHONY: configure-stage4-target-libffi maybe-configure-stage4-target-libffi
+maybe-configure-stage4-target-libffi:
+@if target-libffi-bootstrap
+maybe-configure-stage4-target-libffi: configure-stage4-target-libffi
+configure-stage4-target-libffi:
+	@[ $(current_stage) = stage4 ] || $(MAKE) stage4-start
+	@$(SHELL) $(srcdir)/mkinstalldirs $(TARGET_SUBDIR)/libffi
+	@r=`${PWD_COMMAND}`; export r; \
+	s=`cd $(srcdir); ${PWD_COMMAND}`; export s; \
+	TFLAGS="$(STAGE4_TFLAGS)"; \
+	echo "Checking multilib configuration for libffi..."; \
+	$(CC_FOR_TARGET) --print-multi-lib > $(TARGET_SUBDIR)/libffi/multilib.tmp 2> /dev/null; \
+	if test -r $(TARGET_SUBDIR)/libffi/multilib.out; then \
+	  if cmp -s $(TARGET_SUBDIR)/libffi/multilib.tmp $(TARGET_SUBDIR)/libffi/multilib.out; then \
+	    rm -f $(TARGET_SUBDIR)/libffi/multilib.tmp; \
+	  else \
+	    rm -f $(TARGET_SUBDIR)/libffi/Makefile; \
+	    mv $(TARGET_SUBDIR)/libffi/multilib.tmp $(TARGET_SUBDIR)/libffi/multilib.out; \
+	  fi; \
+	else \
+	  mv $(TARGET_SUBDIR)/libffi/multilib.tmp $(TARGET_SUBDIR)/libffi/multilib.out; \
+	fi; \
+	test ! -f $(TARGET_SUBDIR)/libffi/Makefile || exit 0; \
+	$(NORMAL_TARGET_EXPORTS) \
+	 \
+	CFLAGS="$(CFLAGS_FOR_TARGET)"; export CFLAGS; \
+	CXXFLAGS="$(CXXFLAGS_FOR_TARGET)"; export CXXFLAGS; \
+	LIBCFLAGS="$(LIBCFLAGS_FOR_TARGET)"; export LIBCFLAGS;  \
+	echo Configuring stage 4 in $(TARGET_SUBDIR)/libffi; \
+	$(SHELL) $(srcdir)/mkinstalldirs $(TARGET_SUBDIR)/libffi; \
+	cd $(TARGET_SUBDIR)/libffi || exit 1; \
+	case $(srcdir) in \
+	  /* | [A-Za-z]:[\\/]*) topdir=$(srcdir) ;; \
+	  *) topdir=`echo $(TARGET_SUBDIR)/libffi/ | \
+		sed -e 's,\./,,g' -e 's,[^/]*/,../,g' `$(srcdir) ;; \
+	esac; \
+	module_srcdir=libffi; \
+	$(SHELL) $$s/$$module_srcdir/configure \
+	  --srcdir=$${topdir}/$$module_srcdir \
+	  $(TARGET_CONFIGARGS) --build=${build_alias} --host=${target_alias} \
+	  --target=${target_alias} \
+	  --with-build-libsubdir=$(HOST_SUBDIR) \
+	  $(STAGE4_CONFIGURE_FLAGS)
+@endif target-libffi-bootstrap
+
+.PHONY: configure-stageprofile-target-libffi maybe-configure-stageprofile-target-libffi
+maybe-configure-stageprofile-target-libffi:
+@if target-libffi-bootstrap
+maybe-configure-stageprofile-target-libffi: configure-stageprofile-target-libffi
+configure-stageprofile-target-libffi:
+	@[ $(current_stage) = stageprofile ] || $(MAKE) stageprofile-start
+	@$(SHELL) $(srcdir)/mkinstalldirs $(TARGET_SUBDIR)/libffi
+	@r=`${PWD_COMMAND}`; export r; \
+	s=`cd $(srcdir); ${PWD_COMMAND}`; export s; \
+	TFLAGS="$(STAGEprofile_TFLAGS)"; \
+	echo "Checking multilib configuration for libffi..."; \
+	$(CC_FOR_TARGET) --print-multi-lib > $(TARGET_SUBDIR)/libffi/multilib.tmp 2> /dev/null; \
+	if test -r $(TARGET_SUBDIR)/libffi/multilib.out; then \
+	  if cmp -s $(TARGET_SUBDIR)/libffi/multilib.tmp $(TARGET_SUBDIR)/libffi/multilib.out; then \
+	    rm -f $(TARGET_SUBDIR)/libffi/multilib.tmp; \
+	  else \
+	    rm -f $(TARGET_SUBDIR)/libffi/Makefile; \
+	    mv $(TARGET_SUBDIR)/libffi/multilib.tmp $(TARGET_SUBDIR)/libffi/multilib.out; \
+	  fi; \
+	else \
+	  mv $(TARGET_SUBDIR)/libffi/multilib.tmp $(TARGET_SUBDIR)/libffi/multilib.out; \
+	fi; \
+	test ! -f $(TARGET_SUBDIR)/libffi/Makefile || exit 0; \
+	$(NORMAL_TARGET_EXPORTS) \
+	 \
+	CFLAGS="$(CFLAGS_FOR_TARGET)"; export CFLAGS; \
+	CXXFLAGS="$(CXXFLAGS_FOR_TARGET)"; export CXXFLAGS; \
+	LIBCFLAGS="$(LIBCFLAGS_FOR_TARGET)"; export LIBCFLAGS;  \
+	echo Configuring stage profile in $(TARGET_SUBDIR)/libffi; \
+	$(SHELL) $(srcdir)/mkinstalldirs $(TARGET_SUBDIR)/libffi; \
+	cd $(TARGET_SUBDIR)/libffi || exit 1; \
+	case $(srcdir) in \
+	  /* | [A-Za-z]:[\\/]*) topdir=$(srcdir) ;; \
+	  *) topdir=`echo $(TARGET_SUBDIR)/libffi/ | \
+		sed -e 's,\./,,g' -e 's,[^/]*/,../,g' `$(srcdir) ;; \
+	esac; \
+	module_srcdir=libffi; \
+	$(SHELL) $$s/$$module_srcdir/configure \
+	  --srcdir=$${topdir}/$$module_srcdir \
+	  $(TARGET_CONFIGARGS) --build=${build_alias} --host=${target_alias} \
+	  --target=${target_alias} \
+	  --with-build-libsubdir=$(HOST_SUBDIR) \
+	  $(STAGEprofile_CONFIGURE_FLAGS)
+@endif target-libffi-bootstrap
+
+.PHONY: configure-stagetrain-target-libffi maybe-configure-stagetrain-target-libffi
+maybe-configure-stagetrain-target-libffi:
+@if target-libffi-bootstrap
+maybe-configure-stagetrain-target-libffi: configure-stagetrain-target-libffi
+configure-stagetrain-target-libffi:
+	@[ $(current_stage) = stagetrain ] || $(MAKE) stagetrain-start
+	@$(SHELL) $(srcdir)/mkinstalldirs $(TARGET_SUBDIR)/libffi
+	@r=`${PWD_COMMAND}`; export r; \
+	s=`cd $(srcdir); ${PWD_COMMAND}`; export s; \
+	TFLAGS="$(STAGEtrain_TFLAGS)"; \
+	echo "Checking multilib configuration for libffi..."; \
+	$(CC_FOR_TARGET) --print-multi-lib > $(TARGET_SUBDIR)/libffi/multilib.tmp 2> /dev/null; \
+	if test -r $(TARGET_SUBDIR)/libffi/multilib.out; then \
+	  if cmp -s $(TARGET_SUBDIR)/libffi/multilib.tmp $(TARGET_SUBDIR)/libffi/multilib.out; then \
+	    rm -f $(TARGET_SUBDIR)/libffi/multilib.tmp; \
+	  else \
+	    rm -f $(TARGET_SUBDIR)/libffi/Makefile; \
+	    mv $(TARGET_SUBDIR)/libffi/multilib.tmp $(TARGET_SUBDIR)/libffi/multilib.out; \
+	  fi; \
+	else \
+	  mv $(TARGET_SUBDIR)/libffi/multilib.tmp $(TARGET_SUBDIR)/libffi/multilib.out; \
+	fi; \
+	test ! -f $(TARGET_SUBDIR)/libffi/Makefile || exit 0; \
+	$(NORMAL_TARGET_EXPORTS) \
+	 \
+	CFLAGS="$(CFLAGS_FOR_TARGET)"; export CFLAGS; \
+	CXXFLAGS="$(CXXFLAGS_FOR_TARGET)"; export CXXFLAGS; \
+	LIBCFLAGS="$(LIBCFLAGS_FOR_TARGET)"; export LIBCFLAGS;  \
+	echo Configuring stage train in $(TARGET_SUBDIR)/libffi; \
+	$(SHELL) $(srcdir)/mkinstalldirs $(TARGET_SUBDIR)/libffi; \
+	cd $(TARGET_SUBDIR)/libffi || exit 1; \
+	case $(srcdir) in \
+	  /* | [A-Za-z]:[\\/]*) topdir=$(srcdir) ;; \
+	  *) topdir=`echo $(TARGET_SUBDIR)/libffi/ | \
+		sed -e 's,\./,,g' -e 's,[^/]*/,../,g' `$(srcdir) ;; \
+	esac; \
+	module_srcdir=libffi; \
+	$(SHELL) $$s/$$module_srcdir/configure \
+	  --srcdir=$${topdir}/$$module_srcdir \
+	  $(TARGET_CONFIGARGS) --build=${build_alias} --host=${target_alias} \
+	  --target=${target_alias} \
+	  --with-build-libsubdir=$(HOST_SUBDIR) \
+	  $(STAGEtrain_CONFIGURE_FLAGS)
+@endif target-libffi-bootstrap
+
+.PHONY: configure-stagefeedback-target-libffi maybe-configure-stagefeedback-target-libffi
+maybe-configure-stagefeedback-target-libffi:
+@if target-libffi-bootstrap
+maybe-configure-stagefeedback-target-libffi: configure-stagefeedback-target-libffi
+configure-stagefeedback-target-libffi:
+	@[ $(current_stage) = stagefeedback ] || $(MAKE) stagefeedback-start
+	@$(SHELL) $(srcdir)/mkinstalldirs $(TARGET_SUBDIR)/libffi
+	@r=`${PWD_COMMAND}`; export r; \
+	s=`cd $(srcdir); ${PWD_COMMAND}`; export s; \
+	TFLAGS="$(STAGEfeedback_TFLAGS)"; \
+	echo "Checking multilib configuration for libffi..."; \
+	$(CC_FOR_TARGET) --print-multi-lib > $(TARGET_SUBDIR)/libffi/multilib.tmp 2> /dev/null; \
+	if test -r $(TARGET_SUBDIR)/libffi/multilib.out; then \
+	  if cmp -s $(TARGET_SUBDIR)/libffi/multilib.tmp $(TARGET_SUBDIR)/libffi/multilib.out; then \
+	    rm -f $(TARGET_SUBDIR)/libffi/multilib.tmp; \
+	  else \
+	    rm -f $(TARGET_SUBDIR)/libffi/Makefile; \
+	    mv $(TARGET_SUBDIR)/libffi/multilib.tmp $(TARGET_SUBDIR)/libffi/multilib.out; \
+	  fi; \
+	else \
+	  mv $(TARGET_SUBDIR)/libffi/multilib.tmp $(TARGET_SUBDIR)/libffi/multilib.out; \
+	fi; \
+	test ! -f $(TARGET_SUBDIR)/libffi/Makefile || exit 0; \
+	$(NORMAL_TARGET_EXPORTS) \
+	 \
+	CFLAGS="$(CFLAGS_FOR_TARGET)"; export CFLAGS; \
+	CXXFLAGS="$(CXXFLAGS_FOR_TARGET)"; export CXXFLAGS; \
+	LIBCFLAGS="$(LIBCFLAGS_FOR_TARGET)"; export LIBCFLAGS;  \
+	echo Configuring stage feedback in $(TARGET_SUBDIR)/libffi; \
+	$(SHELL) $(srcdir)/mkinstalldirs $(TARGET_SUBDIR)/libffi; \
+	cd $(TARGET_SUBDIR)/libffi || exit 1; \
+	case $(srcdir) in \
+	  /* | [A-Za-z]:[\\/]*) topdir=$(srcdir) ;; \
+	  *) topdir=`echo $(TARGET_SUBDIR)/libffi/ | \
+		sed -e 's,\./,,g' -e 's,[^/]*/,../,g' `$(srcdir) ;; \
+	esac; \
+	module_srcdir=libffi; \
+	$(SHELL) $$s/$$module_srcdir/configure \
+	  --srcdir=$${topdir}/$$module_srcdir \
+	  $(TARGET_CONFIGARGS) --build=${build_alias} --host=${target_alias} \
+	  --target=${target_alias} \
+	  --with-build-libsubdir=$(HOST_SUBDIR) \
+	  $(STAGEfeedback_CONFIGURE_FLAGS)
+@endif target-libffi-bootstrap
+
+.PHONY: configure-stageautoprofile-target-libffi maybe-configure-stageautoprofile-target-libffi
+maybe-configure-stageautoprofile-target-libffi:
+@if target-libffi-bootstrap
+maybe-configure-stageautoprofile-target-libffi: configure-stageautoprofile-target-libffi
+configure-stageautoprofile-target-libffi:
+	@[ $(current_stage) = stageautoprofile ] || $(MAKE) stageautoprofile-start
+	@$(SHELL) $(srcdir)/mkinstalldirs $(TARGET_SUBDIR)/libffi
+	@r=`${PWD_COMMAND}`; export r; \
+	s=`cd $(srcdir); ${PWD_COMMAND}`; export s; \
+	TFLAGS="$(STAGEautoprofile_TFLAGS)"; \
+	echo "Checking multilib configuration for libffi..."; \
+	$(CC_FOR_TARGET) --print-multi-lib > $(TARGET_SUBDIR)/libffi/multilib.tmp 2> /dev/null; \
+	if test -r $(TARGET_SUBDIR)/libffi/multilib.out; then \
+	  if cmp -s $(TARGET_SUBDIR)/libffi/multilib.tmp $(TARGET_SUBDIR)/libffi/multilib.out; then \
+	    rm -f $(TARGET_SUBDIR)/libffi/multilib.tmp; \
+	  else \
+	    rm -f $(TARGET_SUBDIR)/libffi/Makefile; \
+	    mv $(TARGET_SUBDIR)/libffi/multilib.tmp $(TARGET_SUBDIR)/libffi/multilib.out; \
+	  fi; \
+	else \
+	  mv $(TARGET_SUBDIR)/libffi/multilib.tmp $(TARGET_SUBDIR)/libffi/multilib.out; \
+	fi; \
+	test ! -f $(TARGET_SUBDIR)/libffi/Makefile || exit 0; \
+	$(NORMAL_TARGET_EXPORTS) \
+	 \
+	CFLAGS="$(CFLAGS_FOR_TARGET)"; export CFLAGS; \
+	CXXFLAGS="$(CXXFLAGS_FOR_TARGET)"; export CXXFLAGS; \
+	LIBCFLAGS="$(LIBCFLAGS_FOR_TARGET)"; export LIBCFLAGS;  \
+	echo Configuring stage autoprofile in $(TARGET_SUBDIR)/libffi; \
+	$(SHELL) $(srcdir)/mkinstalldirs $(TARGET_SUBDIR)/libffi; \
+	cd $(TARGET_SUBDIR)/libffi || exit 1; \
+	case $(srcdir) in \
+	  /* | [A-Za-z]:[\\/]*) topdir=$(srcdir) ;; \
+	  *) topdir=`echo $(TARGET_SUBDIR)/libffi/ | \
+		sed -e 's,\./,,g' -e 's,[^/]*/,../,g' `$(srcdir) ;; \
+	esac; \
+	module_srcdir=libffi; \
+	$(SHELL) $$s/$$module_srcdir/configure \
+	  --srcdir=$${topdir}/$$module_srcdir \
+	  $(TARGET_CONFIGARGS) --build=${build_alias} --host=${target_alias} \
+	  --target=${target_alias} \
+	  --with-build-libsubdir=$(HOST_SUBDIR) \
+	  $(STAGEautoprofile_CONFIGURE_FLAGS)
+@endif target-libffi-bootstrap
+
+.PHONY: configure-stageautofeedback-target-libffi maybe-configure-stageautofeedback-target-libffi
+maybe-configure-stageautofeedback-target-libffi:
+@if target-libffi-bootstrap
+maybe-configure-stageautofeedback-target-libffi: configure-stageautofeedback-target-libffi
+configure-stageautofeedback-target-libffi:
+	@[ $(current_stage) = stageautofeedback ] || $(MAKE) stageautofeedback-start
+	@$(SHELL) $(srcdir)/mkinstalldirs $(TARGET_SUBDIR)/libffi
+	@r=`${PWD_COMMAND}`; export r; \
+	s=`cd $(srcdir); ${PWD_COMMAND}`; export s; \
+	TFLAGS="$(STAGEautofeedback_TFLAGS)"; \
+	echo "Checking multilib configuration for libffi..."; \
+	$(CC_FOR_TARGET) --print-multi-lib > $(TARGET_SUBDIR)/libffi/multilib.tmp 2> /dev/null; \
+	if test -r $(TARGET_SUBDIR)/libffi/multilib.out; then \
+	  if cmp -s $(TARGET_SUBDIR)/libffi/multilib.tmp $(TARGET_SUBDIR)/libffi/multilib.out; then \
+	    rm -f $(TARGET_SUBDIR)/libffi/multilib.tmp; \
+	  else \
+	    rm -f $(TARGET_SUBDIR)/libffi/Makefile; \
+	    mv $(TARGET_SUBDIR)/libffi/multilib.tmp $(TARGET_SUBDIR)/libffi/multilib.out; \
+	  fi; \
+	else \
+	  mv $(TARGET_SUBDIR)/libffi/multilib.tmp $(TARGET_SUBDIR)/libffi/multilib.out; \
+	fi; \
+	test ! -f $(TARGET_SUBDIR)/libffi/Makefile || exit 0; \
+	$(NORMAL_TARGET_EXPORTS) \
+	 \
+	CFLAGS="$(CFLAGS_FOR_TARGET)"; export CFLAGS; \
+	CXXFLAGS="$(CXXFLAGS_FOR_TARGET)"; export CXXFLAGS; \
+	LIBCFLAGS="$(LIBCFLAGS_FOR_TARGET)"; export LIBCFLAGS;  \
+	echo Configuring stage autofeedback in $(TARGET_SUBDIR)/libffi; \
+	$(SHELL) $(srcdir)/mkinstalldirs $(TARGET_SUBDIR)/libffi; \
+	cd $(TARGET_SUBDIR)/libffi || exit 1; \
+	case $(srcdir) in \
+	  /* | [A-Za-z]:[\\/]*) topdir=$(srcdir) ;; \
+	  *) topdir=`echo $(TARGET_SUBDIR)/libffi/ | \
+		sed -e 's,\./,,g' -e 's,[^/]*/,../,g' `$(srcdir) ;; \
+	esac; \
+	module_srcdir=libffi; \
+	$(SHELL) $$s/$$module_srcdir/configure \
+	  --srcdir=$${topdir}/$$module_srcdir \
+	  $(TARGET_CONFIGARGS) --build=${build_alias} --host=${target_alias} \
+	  --target=${target_alias} \
+	  --with-build-libsubdir=$(HOST_SUBDIR) \
+	  $(STAGEautofeedback_CONFIGURE_FLAGS)
+@endif target-libffi-bootstrap
+
+
+
 
 
 .PHONY: all-target-libffi maybe-all-target-libffi
@@ -49023,7 +49430,6 @@ all-target-libffi: stage_current
 TARGET-target-libffi=all
 maybe-all-target-libffi: all-target-libffi
 all-target-libffi: configure-target-libffi
-	@: $(MAKE); $(unstage)
 	@r=`${PWD_COMMAND}`; export r; \
 	s=`cd $(srcdir); ${PWD_COMMAND}`; export s; \
 	$(NORMAL_TARGET_EXPORTS)  \
@@ -49034,6 +49440,387 @@ all-target-libffi: configure-target-libffi
 
 
 
+.PHONY: all-stage1-target-libffi maybe-all-stage1-target-libffi
+.PHONY: clean-stage1-target-libffi maybe-clean-stage1-target-libffi
+maybe-all-stage1-target-libffi:
+maybe-clean-stage1-target-libffi:
+@if target-libffi-bootstrap
+maybe-all-stage1-target-libffi: all-stage1-target-libffi
+all-stage1: all-stage1-target-libffi
+TARGET-stage1-target-libffi = $(TARGET-target-libffi)
+all-stage1-target-libffi: configure-stage1-target-libffi
+	@[ $(current_stage) = stage1 ] || $(MAKE) stage1-start
+	@r=`${PWD_COMMAND}`; export r; \
+	s=`cd $(srcdir); ${PWD_COMMAND}`; export s; \
+	TFLAGS="$(STAGE1_TFLAGS)"; \
+	$(NORMAL_TARGET_EXPORTS)  \
+	cd $(TARGET_SUBDIR)/libffi && \
+	 \
+	$(MAKE) $(BASE_FLAGS_TO_PASS) \
+		CFLAGS="$(CFLAGS_FOR_TARGET)" \
+		CXXFLAGS="$(CXXFLAGS_FOR_TARGET)" \
+		LIBCFLAGS="$(LIBCFLAGS_FOR_TARGET)" \
+		CFLAGS_FOR_TARGET="$(CFLAGS_FOR_TARGET)" \
+		CXXFLAGS_FOR_TARGET="$(CXXFLAGS_FOR_TARGET)" \
+		LIBCFLAGS_FOR_TARGET="$(LIBCFLAGS_FOR_TARGET)" \
+		$(EXTRA_TARGET_FLAGS)  \
+		  \
+		TFLAGS="$(STAGE1_TFLAGS)"  \
+		$(TARGET-stage1-target-libffi)
+
+maybe-clean-stage1-target-libffi: clean-stage1-target-libffi
+clean-stage1: clean-stage1-target-libffi
+clean-stage1-target-libffi:
+	@if [ $(current_stage) = stage1 ]; then \
+	  [ -f $(TARGET_SUBDIR)/libffi/Makefile ] || exit 0; \
+	else \
+	  [ -f $(TARGET_SUBDIR)/stage1-libffi/Makefile ] || exit 0; \
+	  $(MAKE) stage1-start; \
+	fi; \
+	cd $(TARGET_SUBDIR)/libffi && \
+	$(MAKE) $(EXTRA_TARGET_FLAGS)  \
+	  clean
+@endif target-libffi-bootstrap
+
+
+.PHONY: all-stage2-target-libffi maybe-all-stage2-target-libffi
+.PHONY: clean-stage2-target-libffi maybe-clean-stage2-target-libffi
+maybe-all-stage2-target-libffi:
+maybe-clean-stage2-target-libffi:
+@if target-libffi-bootstrap
+maybe-all-stage2-target-libffi: all-stage2-target-libffi
+all-stage2: all-stage2-target-libffi
+TARGET-stage2-target-libffi = $(TARGET-target-libffi)
+all-stage2-target-libffi: configure-stage2-target-libffi
+	@[ $(current_stage) = stage2 ] || $(MAKE) stage2-start
+	@r=`${PWD_COMMAND}`; export r; \
+	s=`cd $(srcdir); ${PWD_COMMAND}`; export s; \
+	TFLAGS="$(STAGE2_TFLAGS)"; \
+	$(NORMAL_TARGET_EXPORTS) \
+	  \
+	cd $(TARGET_SUBDIR)/libffi && \
+	 \
+	$(MAKE) $(BASE_FLAGS_TO_PASS) \
+		CFLAGS="$(CFLAGS_FOR_TARGET)" \
+		CXXFLAGS="$(CXXFLAGS_FOR_TARGET)" \
+		LIBCFLAGS="$(LIBCFLAGS_FOR_TARGET)" \
+		CFLAGS_FOR_TARGET="$(CFLAGS_FOR_TARGET)" \
+		CXXFLAGS_FOR_TARGET="$(CXXFLAGS_FOR_TARGET)" \
+		LIBCFLAGS_FOR_TARGET="$(LIBCFLAGS_FOR_TARGET)" \
+		$(EXTRA_TARGET_FLAGS)   \
+		TFLAGS="$(STAGE2_TFLAGS)"  \
+		$(TARGET-stage2-target-libffi)
+
+maybe-clean-stage2-target-libffi: clean-stage2-target-libffi
+clean-stage2: clean-stage2-target-libffi
+clean-stage2-target-libffi:
+	@if [ $(current_stage) = stage2 ]; then \
+	  [ -f $(TARGET_SUBDIR)/libffi/Makefile ] || exit 0; \
+	else \
+	  [ -f $(TARGET_SUBDIR)/stage2-libffi/Makefile ] || exit 0; \
+	  $(MAKE) stage2-start; \
+	fi; \
+	cd $(TARGET_SUBDIR)/libffi && \
+	$(MAKE) $(EXTRA_TARGET_FLAGS)   clean
+@endif target-libffi-bootstrap
+
+
+.PHONY: all-stage3-target-libffi maybe-all-stage3-target-libffi
+.PHONY: clean-stage3-target-libffi maybe-clean-stage3-target-libffi
+maybe-all-stage3-target-libffi:
+maybe-clean-stage3-target-libffi:
+@if target-libffi-bootstrap
+maybe-all-stage3-target-libffi: all-stage3-target-libffi
+all-stage3: all-stage3-target-libffi
+TARGET-stage3-target-libffi = $(TARGET-target-libffi)
+all-stage3-target-libffi: configure-stage3-target-libffi
+	@[ $(current_stage) = stage3 ] || $(MAKE) stage3-start
+	@r=`${PWD_COMMAND}`; export r; \
+	s=`cd $(srcdir); ${PWD_COMMAND}`; export s; \
+	TFLAGS="$(STAGE3_TFLAGS)"; \
+	$(NORMAL_TARGET_EXPORTS) \
+	  \
+	cd $(TARGET_SUBDIR)/libffi && \
+	 \
+	$(MAKE) $(BASE_FLAGS_TO_PASS) \
+		CFLAGS="$(CFLAGS_FOR_TARGET)" \
+		CXXFLAGS="$(CXXFLAGS_FOR_TARGET)" \
+		LIBCFLAGS="$(LIBCFLAGS_FOR_TARGET)" \
+		CFLAGS_FOR_TARGET="$(CFLAGS_FOR_TARGET)" \
+		CXXFLAGS_FOR_TARGET="$(CXXFLAGS_FOR_TARGET)" \
+		LIBCFLAGS_FOR_TARGET="$(LIBCFLAGS_FOR_TARGET)" \
+		$(EXTRA_TARGET_FLAGS)   \
+		TFLAGS="$(STAGE3_TFLAGS)"  \
+		$(TARGET-stage3-target-libffi)
+
+maybe-clean-stage3-target-libffi: clean-stage3-target-libffi
+clean-stage3: clean-stage3-target-libffi
+clean-stage3-target-libffi:
+	@if [ $(current_stage) = stage3 ]; then \
+	  [ -f $(TARGET_SUBDIR)/libffi/Makefile ] || exit 0; \
+	else \
+	  [ -f $(TARGET_SUBDIR)/stage3-libffi/Makefile ] || exit 0; \
+	  $(MAKE) stage3-start; \
+	fi; \
+	cd $(TARGET_SUBDIR)/libffi && \
+	$(MAKE) $(EXTRA_TARGET_FLAGS)   clean
+@endif target-libffi-bootstrap
+
+
+.PHONY: all-stage4-target-libffi maybe-all-stage4-target-libffi
+.PHONY: clean-stage4-target-libffi maybe-clean-stage4-target-libffi
+maybe-all-stage4-target-libffi:
+maybe-clean-stage4-target-libffi:
+@if target-libffi-bootstrap
+maybe-all-stage4-target-libffi: all-stage4-target-libffi
+all-stage4: all-stage4-target-libffi
+TARGET-stage4-target-libffi = $(TARGET-target-libffi)
+all-stage4-target-libffi: configure-stage4-target-libffi
+	@[ $(current_stage) = stage4 ] || $(MAKE) stage4-start
+	@r=`${PWD_COMMAND}`; export r; \
+	s=`cd $(srcdir); ${PWD_COMMAND}`; export s; \
+	TFLAGS="$(STAGE4_TFLAGS)"; \
+	$(NORMAL_TARGET_EXPORTS) \
+	  \
+	cd $(TARGET_SUBDIR)/libffi && \
+	 \
+	$(MAKE) $(BASE_FLAGS_TO_PASS) \
+		CFLAGS="$(CFLAGS_FOR_TARGET)" \
+		CXXFLAGS="$(CXXFLAGS_FOR_TARGET)" \
+		LIBCFLAGS="$(LIBCFLAGS_FOR_TARGET)" \
+		CFLAGS_FOR_TARGET="$(CFLAGS_FOR_TARGET)" \
+		CXXFLAGS_FOR_TARGET="$(CXXFLAGS_FOR_TARGET)" \
+		LIBCFLAGS_FOR_TARGET="$(LIBCFLAGS_FOR_TARGET)" \
+		$(EXTRA_TARGET_FLAGS)   \
+		TFLAGS="$(STAGE4_TFLAGS)"  \
+		$(TARGET-stage4-target-libffi)
+
+maybe-clean-stage4-target-libffi: clean-stage4-target-libffi
+clean-stage4: clean-stage4-target-libffi
+clean-stage4-target-libffi:
+	@if [ $(current_stage) = stage4 ]; then \
+	  [ -f $(TARGET_SUBDIR)/libffi/Makefile ] || exit 0; \
+	else \
+	  [ -f $(TARGET_SUBDIR)/stage4-libffi/Makefile ] || exit 0; \
+	  $(MAKE) stage4-start; \
+	fi; \
+	cd $(TARGET_SUBDIR)/libffi && \
+	$(MAKE) $(EXTRA_TARGET_FLAGS)   clean
+@endif target-libffi-bootstrap
+
+
+.PHONY: all-stageprofile-target-libffi maybe-all-stageprofile-target-libffi
+.PHONY: clean-stageprofile-target-libffi maybe-clean-stageprofile-target-libffi
+maybe-all-stageprofile-target-libffi:
+maybe-clean-stageprofile-target-libffi:
+@if target-libffi-bootstrap
+maybe-all-stageprofile-target-libffi: all-stageprofile-target-libffi
+all-stageprofile: all-stageprofile-target-libffi
+TARGET-stageprofile-target-libffi = $(TARGET-target-libffi)
+all-stageprofile-target-libffi: configure-stageprofile-target-libffi
+	@[ $(current_stage) = stageprofile ] || $(MAKE) stageprofile-start
+	@r=`${PWD_COMMAND}`; export r; \
+	s=`cd $(srcdir); ${PWD_COMMAND}`; export s; \
+	TFLAGS="$(STAGEprofile_TFLAGS)"; \
+	$(NORMAL_TARGET_EXPORTS) \
+	  \
+	cd $(TARGET_SUBDIR)/libffi && \
+	 \
+	$(MAKE) $(BASE_FLAGS_TO_PASS) \
+		CFLAGS="$(CFLAGS_FOR_TARGET)" \
+		CXXFLAGS="$(CXXFLAGS_FOR_TARGET)" \
+		LIBCFLAGS="$(LIBCFLAGS_FOR_TARGET)" \
+		CFLAGS_FOR_TARGET="$(CFLAGS_FOR_TARGET)" \
+		CXXFLAGS_FOR_TARGET="$(CXXFLAGS_FOR_TARGET)" \
+		LIBCFLAGS_FOR_TARGET="$(LIBCFLAGS_FOR_TARGET)" \
+		$(EXTRA_TARGET_FLAGS)   \
+		TFLAGS="$(STAGEprofile_TFLAGS)"  \
+		$(TARGET-stageprofile-target-libffi)
+
+maybe-clean-stageprofile-target-libffi: clean-stageprofile-target-libffi
+clean-stageprofile: clean-stageprofile-target-libffi
+clean-stageprofile-target-libffi:
+	@if [ $(current_stage) = stageprofile ]; then \
+	  [ -f $(TARGET_SUBDIR)/libffi/Makefile ] || exit 0; \
+	else \
+	  [ -f $(TARGET_SUBDIR)/stageprofile-libffi/Makefile ] || exit 0; \
+	  $(MAKE) stageprofile-start; \
+	fi; \
+	cd $(TARGET_SUBDIR)/libffi && \
+	$(MAKE) $(EXTRA_TARGET_FLAGS)   clean
+@endif target-libffi-bootstrap
+
+
+.PHONY: all-stagetrain-target-libffi maybe-all-stagetrain-target-libffi
+.PHONY: clean-stagetrain-target-libffi maybe-clean-stagetrain-target-libffi
+maybe-all-stagetrain-target-libffi:
+maybe-clean-stagetrain-target-libffi:
+@if target-libffi-bootstrap
+maybe-all-stagetrain-target-libffi: all-stagetrain-target-libffi
+all-stagetrain: all-stagetrain-target-libffi
+TARGET-stagetrain-target-libffi = $(TARGET-target-libffi)
+all-stagetrain-target-libffi: configure-stagetrain-target-libffi
+	@[ $(current_stage) = stagetrain ] || $(MAKE) stagetrain-start
+	@r=`${PWD_COMMAND}`; export r; \
+	s=`cd $(srcdir); ${PWD_COMMAND}`; export s; \
+	TFLAGS="$(STAGEtrain_TFLAGS)"; \
+	$(NORMAL_TARGET_EXPORTS) \
+	  \
+	cd $(TARGET_SUBDIR)/libffi && \
+	 \
+	$(MAKE) $(BASE_FLAGS_TO_PASS) \
+		CFLAGS="$(CFLAGS_FOR_TARGET)" \
+		CXXFLAGS="$(CXXFLAGS_FOR_TARGET)" \
+		LIBCFLAGS="$(LIBCFLAGS_FOR_TARGET)" \
+		CFLAGS_FOR_TARGET="$(CFLAGS_FOR_TARGET)" \
+		CXXFLAGS_FOR_TARGET="$(CXXFLAGS_FOR_TARGET)" \
+		LIBCFLAGS_FOR_TARGET="$(LIBCFLAGS_FOR_TARGET)" \
+		$(EXTRA_TARGET_FLAGS)   \
+		TFLAGS="$(STAGEtrain_TFLAGS)"  \
+		$(TARGET-stagetrain-target-libffi)
+
+maybe-clean-stagetrain-target-libffi: clean-stagetrain-target-libffi
+clean-stagetrain: clean-stagetrain-target-libffi
+clean-stagetrain-target-libffi:
+	@if [ $(current_stage) = stagetrain ]; then \
+	  [ -f $(TARGET_SUBDIR)/libffi/Makefile ] || exit 0; \
+	else \
+	  [ -f $(TARGET_SUBDIR)/stagetrain-libffi/Makefile ] || exit 0; \
+	  $(MAKE) stagetrain-start; \
+	fi; \
+	cd $(TARGET_SUBDIR)/libffi && \
+	$(MAKE) $(EXTRA_TARGET_FLAGS)   clean
+@endif target-libffi-bootstrap
+
+
+.PHONY: all-stagefeedback-target-libffi maybe-all-stagefeedback-target-libffi
+.PHONY: clean-stagefeedback-target-libffi maybe-clean-stagefeedback-target-libffi
+maybe-all-stagefeedback-target-libffi:
+maybe-clean-stagefeedback-target-libffi:
+@if target-libffi-bootstrap
+maybe-all-stagefeedback-target-libffi: all-stagefeedback-target-libffi
+all-stagefeedback: all-stagefeedback-target-libffi
+TARGET-stagefeedback-target-libffi = $(TARGET-target-libffi)
+all-stagefeedback-target-libffi: configure-stagefeedback-target-libffi
+	@[ $(current_stage) = stagefeedback ] || $(MAKE) stagefeedback-start
+	@r=`${PWD_COMMAND}`; export r; \
+	s=`cd $(srcdir); ${PWD_COMMAND}`; export s; \
+	TFLAGS="$(STAGEfeedback_TFLAGS)"; \
+	$(NORMAL_TARGET_EXPORTS) \
+	  \
+	cd $(TARGET_SUBDIR)/libffi && \
+	 \
+	$(MAKE) $(BASE_FLAGS_TO_PASS) \
+		CFLAGS="$(CFLAGS_FOR_TARGET)" \
+		CXXFLAGS="$(CXXFLAGS_FOR_TARGET)" \
+		LIBCFLAGS="$(LIBCFLAGS_FOR_TARGET)" \
+		CFLAGS_FOR_TARGET="$(CFLAGS_FOR_TARGET)" \
+		CXXFLAGS_FOR_TARGET="$(CXXFLAGS_FOR_TARGET)" \
+		LIBCFLAGS_FOR_TARGET="$(LIBCFLAGS_FOR_TARGET)" \
+		$(EXTRA_TARGET_FLAGS)   \
+		TFLAGS="$(STAGEfeedback_TFLAGS)"  \
+		$(TARGET-stagefeedback-target-libffi)
+
+maybe-clean-stagefeedback-target-libffi: clean-stagefeedback-target-libffi
+clean-stagefeedback: clean-stagefeedback-target-libffi
+clean-stagefeedback-target-libffi:
+	@if [ $(current_stage) = stagefeedback ]; then \
+	  [ -f $(TARGET_SUBDIR)/libffi/Makefile ] || exit 0; \
+	else \
+	  [ -f $(TARGET_SUBDIR)/stagefeedback-libffi/Makefile ] || exit 0; \
+	  $(MAKE) stagefeedback-start; \
+	fi; \
+	cd $(TARGET_SUBDIR)/libffi && \
+	$(MAKE) $(EXTRA_TARGET_FLAGS)   clean
+@endif target-libffi-bootstrap
+
+
+.PHONY: all-stageautoprofile-target-libffi maybe-all-stageautoprofile-target-libffi
+.PHONY: clean-stageautoprofile-target-libffi maybe-clean-stageautoprofile-target-libffi
+maybe-all-stageautoprofile-target-libffi:
+maybe-clean-stageautoprofile-target-libffi:
+@if target-libffi-bootstrap
+maybe-all-stageautoprofile-target-libffi: all-stageautoprofile-target-libffi
+all-stageautoprofile: all-stageautoprofile-target-libffi
+TARGET-stageautoprofile-target-libffi = $(TARGET-target-libffi)
+all-stageautoprofile-target-libffi: configure-stageautoprofile-target-libffi
+	@[ $(current_stage) = stageautoprofile ] || $(MAKE) stageautoprofile-start
+	@r=`${PWD_COMMAND}`; export r; \
+	s=`cd $(srcdir); ${PWD_COMMAND}`; export s; \
+	TFLAGS="$(STAGEautoprofile_TFLAGS)"; \
+	$(NORMAL_TARGET_EXPORTS) \
+	  \
+	cd $(TARGET_SUBDIR)/libffi && \
+	$$s/gcc/config/i386/$(AUTO_PROFILE) \
+	$(MAKE) $(BASE_FLAGS_TO_PASS) \
+		CFLAGS="$(CFLAGS_FOR_TARGET)" \
+		CXXFLAGS="$(CXXFLAGS_FOR_TARGET)" \
+		LIBCFLAGS="$(LIBCFLAGS_FOR_TARGET)" \
+		CFLAGS_FOR_TARGET="$(CFLAGS_FOR_TARGET)" \
+		CXXFLAGS_FOR_TARGET="$(CXXFLAGS_FOR_TARGET)" \
+		LIBCFLAGS_FOR_TARGET="$(LIBCFLAGS_FOR_TARGET)" \
+		$(EXTRA_TARGET_FLAGS)   \
+		TFLAGS="$(STAGEautoprofile_TFLAGS)"  \
+		$(TARGET-stageautoprofile-target-libffi)
+
+maybe-clean-stageautoprofile-target-libffi: clean-stageautoprofile-target-libffi
+clean-stageautoprofile: clean-stageautoprofile-target-libffi
+clean-stageautoprofile-target-libffi:
+	@if [ $(current_stage) = stageautoprofile ]; then \
+	  [ -f $(TARGET_SUBDIR)/libffi/Makefile ] || exit 0; \
+	else \
+	  [ -f $(TARGET_SUBDIR)/stageautoprofile-libffi/Makefile ] || exit 0; \
+	  $(MAKE) stageautoprofile-start; \
+	fi; \
+	cd $(TARGET_SUBDIR)/libffi && \
+	$(MAKE) $(EXTRA_TARGET_FLAGS)   clean
+@endif target-libffi-bootstrap
+
+
+.PHONY: all-stageautofeedback-target-libffi maybe-all-stageautofeedback-target-libffi
+.PHONY: clean-stageautofeedback-target-libffi maybe-clean-stageautofeedback-target-libffi
+maybe-all-stageautofeedback-target-libffi:
+maybe-clean-stageautofeedback-target-libffi:
+@if target-libffi-bootstrap
+maybe-all-stageautofeedback-target-libffi: all-stageautofeedback-target-libffi
+all-stageautofeedback: all-stageautofeedback-target-libffi
+TARGET-stageautofeedback-target-libffi = $(TARGET-target-libffi)
+all-stageautofeedback-target-libffi: configure-stageautofeedback-target-libffi
+	@[ $(current_stage) = stageautofeedback ] || $(MAKE) stageautofeedback-start
+	@r=`${PWD_COMMAND}`; export r; \
+	s=`cd $(srcdir); ${PWD_COMMAND}`; export s; \
+	TFLAGS="$(STAGEautofeedback_TFLAGS)"; \
+	$(NORMAL_TARGET_EXPORTS) \
+	  \
+	cd $(TARGET_SUBDIR)/libffi && \
+	 \
+	$(MAKE) $(BASE_FLAGS_TO_PASS) \
+		CFLAGS="$(CFLAGS_FOR_TARGET)" \
+		CXXFLAGS="$(CXXFLAGS_FOR_TARGET)" \
+		LIBCFLAGS="$(LIBCFLAGS_FOR_TARGET)" \
+		CFLAGS_FOR_TARGET="$(CFLAGS_FOR_TARGET)" \
+		CXXFLAGS_FOR_TARGET="$(CXXFLAGS_FOR_TARGET)" \
+		LIBCFLAGS_FOR_TARGET="$(LIBCFLAGS_FOR_TARGET)" \
+		$(EXTRA_TARGET_FLAGS)   \
+		TFLAGS="$(STAGEautofeedback_TFLAGS)" PERF_DATA=perf.data \
+		$(TARGET-stageautofeedback-target-libffi)
+
+maybe-clean-stageautofeedback-target-libffi: clean-stageautofeedback-target-libffi
+clean-stageautofeedback: clean-stageautofeedback-target-libffi
+clean-stageautofeedback-target-libffi:
+	@if [ $(current_stage) = stageautofeedback ]; then \
+	  [ -f $(TARGET_SUBDIR)/libffi/Makefile ] || exit 0; \
+	else \
+	  [ -f $(TARGET_SUBDIR)/stageautofeedback-libffi/Makefile ] || exit 0; \
+	  $(MAKE) stageautofeedback-start; \
+	fi; \
+	cd $(TARGET_SUBDIR)/libffi && \
+	$(MAKE) $(EXTRA_TARGET_FLAGS)   clean
+@endif target-libffi-bootstrap
+
+
+
+
 
 
 .PHONY: check-target-libffi maybe-check-target-libffi
@@ -56251,7 +57038,15 @@ configure-target-libphobos: stage_last
 configure-target-libtermcap: stage_last
 configure-target-winsup: stage_last
 configure-target-libgloss: stage_last
-configure-target-libffi: stage_last
+configure-stage1-target-libffi: maybe-all-stage1-gcc
+configure-stage2-target-libffi: maybe-all-stage2-gcc
+configure-stage3-target-libffi: maybe-all-stage3-gcc
+configure-stage4-target-libffi: maybe-all-stage4-gcc
+configure-stageprofile-target-libffi: maybe-all-stageprofile-gcc
+configure-stagetrain-target-libffi: maybe-all-stagetrain-gcc
+configure-stagefeedback-target-libffi: maybe-all-stagefeedback-gcc
+configure-stageautoprofile-target-libffi: maybe-all-stageautoprofile-gcc
+configure-stageautofeedback-target-libffi: maybe-all-stageautofeedback-gcc
 configure-target-zlib: stage_last
 configure-target-rda: stage_last
 configure-target-libada: stage_last
@@ -57274,10 +58069,28 @@ all-flex: maybe-all-build-texinfo
 all-m4: maybe-all-build-texinfo
 configure-target-fastjar: maybe-configure-target-zlib
 all-target-fastjar: maybe-all-target-zlib
-configure-target-libgo: maybe-configure-target-libffi
 all-target-libgo: maybe-all-target-libbacktrace
-all-target-libgo: maybe-all-target-libffi
 all-target-libgo: maybe-all-target-libatomic
+all-target-libgomp: maybe-all-target-libffi
+all-stage1-target-libgomp: maybe-all-stage1-target-libffi
+all-stage2-target-libgomp: maybe-all-stage2-target-libffi
+all-stage3-target-libgomp: maybe-all-stage3-target-libffi
+all-stage4-target-libgomp: maybe-all-stage4-target-libffi
+all-stageprofile-target-libgomp: maybe-all-stageprofile-target-libffi
+all-stagetrain-target-libgomp: maybe-all-stagetrain-target-libffi
+all-stagefeedback-target-libgomp: maybe-all-stagefeedback-target-libffi
+all-stageautoprofile-target-libgomp: maybe-all-stageautoprofile-target-libffi
+all-stageautofeedback-target-libgomp: maybe-all-stageautofeedback-target-libffi
+configure-target-libgomp: maybe-configure-target-libffi
+configure-stage1-target-libgomp: maybe-configure-stage1-target-libffi
+configure-stage2-target-libgomp: maybe-configure-stage2-target-libffi
+configure-stage3-target-libgomp: maybe-configure-stage3-target-libffi
+configure-stage4-target-libgomp: maybe-configure-stage4-target-libffi
+configure-stageprofile-target-libgomp: maybe-configure-stageprofile-target-libffi
+configure-stagetrain-target-libgomp: maybe-configure-stagetrain-target-libffi
+configure-stagefeedback-target-libgomp: maybe-configure-stagefeedback-target-libffi
+configure-stageautoprofile-target-libgomp: maybe-configure-stageautoprofile-target-libffi
+configure-stageautofeedback-target-libgomp: maybe-configure-stageautofeedback-target-libffi
 configure-target-libphobos: maybe-configure-target-libbacktrace
 configure-target-libphobos: maybe-configure-target-zlib
 all-target-libphobos: maybe-all-target-libbacktrace
@@ -57328,6 +58141,7 @@ install-target-libgfortran: maybe-install-target-libquadmath
 install-target-libgfortran: maybe-install-target-libgcc
 install-target-libphobos: maybe-install-target-libatomic
 install-target-libsanitizer: maybe-install-target-libstdc++-v3
+install-target-libgomp: maybe-install-target-libffi
 install-target-libsanitizer: maybe-install-target-libgcc
 install-target-libvtv: maybe-install-target-libstdc++-v3
 install-target-libvtv: maybe-install-target-libgcc
@@ -57387,7 +58201,9 @@ all-fastjar: maybe-all-libiberty
 all-bison: maybe-all-intl
 all-flex: maybe-all-intl
 all-m4: maybe-all-intl
+configure-target-libgo: maybe-configure-target-libffi
 configure-target-libgo: maybe-all-target-libstdc++-v3
+all-target-libgo: maybe-all-target-libffi
 configure-target-liboffloadmic: maybe-configure-target-libgomp
 all-target-liboffloadmic: maybe-all-target-libgomp
 configure-target-newlib: maybe-all-binutils
@@ -57427,6 +58243,15 @@ configure-stagetrain-target-libvtv: maybe-all-stagetrain-target-libgcc
 configure-stagefeedback-target-libvtv: maybe-all-stagefeedback-target-libgcc
 configure-stageautoprofile-target-libvtv: maybe-all-stageautoprofile-target-libgcc
 configure-stageautofeedback-target-libvtv: maybe-all-stageautofeedback-target-libgcc
+configure-stage1-target-libffi: maybe-all-stage1-target-libgcc
+configure-stage2-target-libffi: maybe-all-stage2-target-libgcc
+configure-stage3-target-libffi: maybe-all-stage3-target-libgcc
+configure-stage4-target-libffi: maybe-all-stage4-target-libgcc
+configure-stageprofile-target-libffi: maybe-all-stageprofile-target-libgcc
+configure-stagetrain-target-libffi: maybe-all-stagetrain-target-libgcc
+configure-stagefeedback-target-libffi: maybe-all-stagefeedback-target-libgcc
+configure-stageautoprofile-target-libffi: maybe-all-stageautoprofile-target-libgcc
+configure-stageautofeedback-target-libffi: maybe-all-stageautofeedback-target-libgcc
 configure-stage1-target-libgomp: maybe-all-stage1-target-libgcc
 configure-stage2-target-libgomp: maybe-all-stage2-target-libgcc
 configure-stage3-target-libgomp: maybe-all-stage3-target-libgcc
diff --git a/configure b/configure
index abd93a990a94..033929b0ab8c 100755
--- a/configure
+++ b/configure
@@ -3513,11 +3513,19 @@ case "${target}" in
   ft32-*-*)
     noconfigdirs="$noconfigdirs target-libffi"
     ;;
+  nvptx-*-*)
+    noconfigdirs="$noconfigdirs target-libffi"
+    ;;
   *-*-lynxos*)
     noconfigdirs="$noconfigdirs target-libffi"
     ;;
 esac
 
+libgomp_deps="target-libffi"
+if echo " ${noconfigdirs} " | grep " target-libffi " > /dev/null 2>&1 ; then
+   libgomp_deps=""
+fi
+
 # Disable the go frontend on systems where it is known to not work. Please keep
 # this in sync with contrib/config-list.mk.
 case "${target}" in
@@ -6588,6 +6596,15 @@ esac
 # $build_configdirs and $target_configdirs.
 # If we have the source for $noconfigdirs entries, add them to $notsupp.
 
+# libgomp depends on libffi.  Remove it from nonsupp if necessary.
+if ! (echo " $noconfigdirs " | grep " target-libgomp " >/dev/null 2>&1); then
+  if echo " $noconfigdirs " | grep " target-libffi " >/dev/null 2>&1; then
+    if test "x${libgomp_deps}" != x; then
+      noconfigdirs=`echo " $noconfigdirs " | sed -e "s/ target-libffi / /"`
+    fi
+  fi
+fi
+
 notsupp=""
 for dir in . $skipdirs $noconfigdirs ; do
   dirname=`echo $dir | sed -e s/target-//g -e s/build-//g`
@@ -7182,6 +7199,9 @@ bootstrap_fixincludes=no
 
 # If we are building libgomp, bootstrap it.
 if echo " ${target_configdirs} " | grep " libgomp " > /dev/null 2>&1 ; then
+   if echo " ${target_configdirs} " | grep " libffi " > /dev/null 2>&1 ; then
+     bootstrap_target_libs=${bootstrap_target_libs}target-libffi,
+   fi
   bootstrap_target_libs=${bootstrap_target_libs}target-libgomp,
 fi
 
diff --git a/configure.ac b/configure.ac
index 9db4fd14aa23..de361880ba7c 100644
--- a/configure.ac
+++ b/configure.ac
@@ -795,11 +795,19 @@ case "${target}" in
   ft32-*-*)
     noconfigdirs="$noconfigdirs target-libffi"
     ;;
+  nvptx-*-*)
+    noconfigdirs="$noconfigdirs target-libffi"
+    ;;
   *-*-lynxos*)
     noconfigdirs="$noconfigdirs target-libffi"
     ;;
 esac
 
+libgomp_deps="target-libffi"
+if echo " ${noconfigdirs} " | grep " target-libffi " > /dev/null 2>&1 ; then
+   libgomp_deps=""
+fi
+
 # Disable the go frontend on systems where it is known to not work. Please keep
 # this in sync with contrib/config-list.mk.
 case "${target}" in
@@ -2178,6 +2186,15 @@ esac
 # $build_configdirs and $target_configdirs.
 # If we have the source for $noconfigdirs entries, add them to $notsupp.
 
+# libgomp depends on libffi.  Remove it from nonsupp if necessary.
+if ! (echo " $noconfigdirs " | grep " target-libgomp " >/dev/null 2>&1); then
+  if echo " $noconfigdirs " | grep " target-libffi " >/dev/null 2>&1; then
+    if test "x${libgomp_deps}" != x; then
+      noconfigdirs=`echo " $noconfigdirs " | sed -e "s/ target-libffi / /"`
+    fi
+  fi
+fi
+
 notsupp=""
 for dir in . $skipdirs $noconfigdirs ; do
   dirname=`echo $dir | sed -e s/target-//g -e s/build-//g`
@@ -2695,6 +2712,9 @@ bootstrap_fixincludes=no
 
 # If we are building libgomp, bootstrap it.
 if echo " ${target_configdirs} " | grep " libgomp " > /dev/null 2>&1 ; then
+   if echo " ${target_configdirs} " | grep " libffi " > /dev/null 2>&1 ; then
+     bootstrap_target_libs=${bootstrap_target_libs}target-libffi,
+   fi
   bootstrap_target_libs=${bootstrap_target_libs}target-libgomp,
 fi
 
diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp
index 78f4042c5e24..e68675163ea9 100644
--- a/gcc/ChangeLog.omp
+++ b/gcc/ChangeLog.omp
@@ -1,3 +1,8 @@
+2017-12-21  Cesar Philippidis  <cesar@codesourcery.com>
+
+	* omp-low.c (install_parm_decl): Don't extract identifiers from
+	artifical decls.
+
 2018-12-21  Gergö Barany  <gergo@codesourcery.com>
 
 	* omp-expand.c (expand_omp_target): Handle if_present flag on
diff --git a/gcc/builtin-types.def b/gcc/builtin-types.def
index e5c9e063c480..9ee86b4957e3 100644
--- a/gcc/builtin-types.def
+++ b/gcc/builtin-types.def
@@ -826,6 +826,10 @@ DEF_FUNCTION_TYPE_VAR_6 (BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_VAR,
 			 BT_VOID, BT_INT, BT_PTR_FN_VOID_PTR, BT_SIZE,
 			 BT_PTR, BT_PTR, BT_PTR)
 
+DEF_FUNCTION_TYPE_VAR_7 (BT_FN_VOID_INT_INT_OMPFN_SIZE_PTR_PTR_PTR_VAR,
+			 BT_VOID, BT_INT, BT_INT, BT_PTR_FN_VOID_PTR, BT_SIZE,
+			 BT_PTR, BT_PTR, BT_PTR)
+
 DEF_FUNCTION_TYPE_VAR_7 (BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_INT_INT_VAR,
 			 BT_VOID, BT_INT, BT_SIZE, BT_PTR, BT_PTR,
 			 BT_PTR, BT_INT, BT_INT)
diff --git a/gcc/config/nvptx/nvptx.c b/gcc/config/nvptx/nvptx.c
index 0b4feec2e9ec..077f6cc145eb 100644
--- a/gcc/config/nvptx/nvptx.c
+++ b/gcc/config/nvptx/nvptx.c
@@ -5432,6 +5432,10 @@ nvptx_expand_cmp_swap (tree exp, rtx target,
 			 NULL_RTX, mode, EXPAND_NORMAL);
   rtx pat;
 
+  /* 'mem' might be a PARM_DECL.  If so, convert it to a register.  */
+  if (!REG_P (mem))
+    mem = copy_to_mode_reg (GET_MODE (mem), mem);
+
   mem = gen_rtx_MEM (mode, mem);
   if (!REG_P (cmp))
     cmp = copy_to_mode_reg (mode, cmp);
diff --git a/gcc/fortran/ChangeLog.omp b/gcc/fortran/ChangeLog.omp
index 00ccd26462c4..3bf4f43730d4 100644
--- a/gcc/fortran/ChangeLog.omp
+++ b/gcc/fortran/ChangeLog.omp
@@ -1,3 +1,8 @@
+2017-12-21  Cesar Philippidis  <cesar@codesourcery.com>
+
+	* types.def: (BF_FN_VOID_INT_INT_OMPFN_SIZE_PTR_PTR_PTR_VAR):
+	Define.
+
 2018-12-21  Gergö Barany  <gergo@codesourcery.com>
 
 	* openmp.c (OACC_HOST_DATA_CLAUSES): Add OMP_CLAUSE_IF and
diff --git a/gcc/fortran/types.def b/gcc/fortran/types.def
index b96e292fc810..5c976338ff6a 100644
--- a/gcc/fortran/types.def
+++ b/gcc/fortran/types.def
@@ -272,3 +272,7 @@ DEF_FUNCTION_TYPE_VAR_7 (BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_INT_INT_VAR,
 DEF_FUNCTION_TYPE_VAR_6 (BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_VAR,
 			  BT_VOID, BT_INT, BT_PTR_FN_VOID_PTR, BT_SIZE,
 			  BT_PTR, BT_PTR, BT_PTR)
+
+DEF_FUNCTION_TYPE_VAR_7 (BT_FN_VOID_INT_INT_OMPFN_SIZE_PTR_PTR_PTR_VAR,
+			  BT_VOID, BT_INT, BT_INT, BT_PTR_FN_VOID_PTR, BT_SIZE,
+			  BT_PTR, BT_PTR, BT_PTR)
diff --git a/gcc/omp-builtins.def b/gcc/omp-builtins.def
index 9961c2874943..99cb8fa336be 100644
--- a/gcc/omp-builtins.def
+++ b/gcc/omp-builtins.def
@@ -38,8 +38,8 @@ DEF_GOACC_BUILTIN (BUILT_IN_GOACC_DATA_END, "GOACC_data_end",
 DEF_GOACC_BUILTIN (BUILT_IN_GOACC_ENTER_EXIT_DATA, "GOACC_enter_exit_data",
 		   BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_INT_INT_VAR,
 		   ATTR_NOTHROW_LIST)
-DEF_GOACC_BUILTIN (BUILT_IN_GOACC_PARALLEL, "GOACC_parallel_keyed",
-		   BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_VAR,
+DEF_GOACC_BUILTIN (BUILT_IN_GOACC_PARALLEL, "GOACC_parallel_keyed_v2",
+		   BT_FN_VOID_INT_INT_OMPFN_SIZE_PTR_PTR_PTR_VAR,
 		   ATTR_NOTHROW_LIST)
 DEF_GOACC_BUILTIN (BUILT_IN_GOACC_UPDATE, "GOACC_update",
 		   BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_INT_INT_VAR,
diff --git a/gcc/omp-expand.c b/gcc/omp-expand.c
index 6dd87689e5ab..c16db789d9f6 100644
--- a/gcc/omp-expand.c
+++ b/gcc/omp-expand.c
@@ -7306,19 +7306,22 @@ expand_omp_target (struct omp_region *region)
   gomp_target *entry_stmt;
   gimple *stmt;
   edge e;
-  bool offloaded, data_region;
+  bool offloaded, data_region, oacc_parallel;
 
   entry_stmt = as_a <gomp_target *> (last_stmt (region->entry));
   new_bb = region->entry;
+  oacc_parallel = false;
 
   offloaded = is_gimple_omp_offloaded (entry_stmt);
   switch (gimple_omp_target_kind (entry_stmt))
     {
+    case GF_OMP_TARGET_KIND_OACC_PARALLEL:
+      oacc_parallel = true;
+      gcc_fallthrough ();
     case GF_OMP_TARGET_KIND_REGION:
     case GF_OMP_TARGET_KIND_UPDATE:
     case GF_OMP_TARGET_KIND_ENTER_DATA:
     case GF_OMP_TARGET_KIND_EXIT_DATA:
-    case GF_OMP_TARGET_KIND_OACC_PARALLEL:
     case GF_OMP_TARGET_KIND_OACC_KERNELS:
     case GF_OMP_TARGET_KIND_OACC_UPDATE:
     case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
@@ -7380,7 +7383,7 @@ expand_omp_target (struct omp_region *region)
 	 .OMP_DATA_I may have been converted into a different local
 	 variable.  In which case, we need to keep the assignment.  */
       tree data_arg = gimple_omp_target_data_arg (entry_stmt);
-      if (data_arg)
+      if (data_arg && !oacc_parallel)
 	{
 	  basic_block entry_succ_bb = single_succ (entry_bb);
 	  gimple_stmt_iterator gsi;
@@ -7742,6 +7745,11 @@ expand_omp_target (struct omp_region *region)
     }
   else
     args.quick_push (device);
+  if (start_ix == BUILT_IN_GOACC_PARALLEL)
+    {
+      tree use_params = oacc_parallel ? integer_one_node : integer_zero_node;
+      args.quick_push (use_params);
+    }
   if (offloaded)
     args.quick_push (build_fold_addr_expr (child_fn));
   args.quick_push (t1);
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index c50b8ff6eec8..0c8acb27614f 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -91,6 +91,7 @@ struct omp_context
   /* Map variables to fields in a structure that allows communication
      between sending and receiving threads.  */
   splay_tree field_map;
+  splay_tree parm_map;
   tree record_type;
   tree sender_decl;
   tree receiver_decl;
@@ -350,6 +351,14 @@ maybe_lookup_decl (const_tree var, omp_context *ctx)
   return n ? *n : NULL_TREE;
 }
 
+static inline tree
+lookup_parm (const_tree var, omp_context *ctx)
+{
+  splay_tree_node n;
+  n = splay_tree_lookup (ctx->parm_map, (splay_tree_key) var);
+  return (tree) n->value;
+}
+
 static inline tree
 lookup_field (tree var, omp_context *ctx)
 {
@@ -531,15 +540,21 @@ build_receiver_ref (tree var, bool by_ref, omp_context *ctx)
 {
   tree x, field = lookup_field (var, ctx);
 
-  /* If the receiver record type was remapped in the child function,
-     remap the field into the new record type.  */
-  x = maybe_lookup_field (field, ctx);
-  if (x != NULL)
-    field = x;
+  if (is_oacc_parallel (ctx))
+    x = lookup_parm (var, ctx);
+  else
+    {
+      /* If the receiver record type was remapped in the child function,
+	 remap the field into the new record type.  */
+      x = maybe_lookup_field (field, ctx);
+      if (x != NULL)
+	field = x;
+
+      x = build_simple_mem_ref (ctx->receiver_decl);
+      TREE_THIS_NOTRAP (x) = 1;
+      x = omp_build_component_ref (x, field);
+    }
 
-  x = build_simple_mem_ref (ctx->receiver_decl);
-  TREE_THIS_NOTRAP (x) = 1;
-  x = omp_build_component_ref (x, field);
   if (by_ref)
     {
       x = build_simple_mem_ref (x);
@@ -676,6 +691,32 @@ build_sender_ref (tree var, omp_context *ctx)
   return build_sender_ref ((splay_tree_key) var, ctx);
 }
 
+static void
+install_parm_decl (tree var, tree type, omp_context *ctx)
+{
+  if (!is_oacc_parallel (ctx))
+    return;
+
+  splay_tree_key key = (splay_tree_key) var;
+  tree decl_name = NULL_TREE, t;
+  location_t loc = UNKNOWN_LOCATION;
+
+  if (DECL_P (var) && !DECL_ARTIFICIAL (var))
+    {
+      decl_name = get_identifier (get_name (var));
+      loc = DECL_SOURCE_LOCATION (var);
+    }
+  t = build_decl (loc, PARM_DECL, decl_name, type);
+  DECL_ARTIFICIAL (t) = 1;
+  DECL_NAMELESS (t) = 1;
+  DECL_ARG_TYPE (t) = type;
+  DECL_CONTEXT (t) = current_function_decl;
+  TREE_USED (t) = 1;
+  TREE_READONLY (t) = 1;
+
+  splay_tree_insert (ctx->parm_map, key, (splay_tree_value) t);
+}
+
 /* Add a new field for VAR inside the structure CTX->SENDER_DECL.  If
    BASE_POINTERS_RESTRICT, declare the field with restrict.  */
 
@@ -775,7 +816,10 @@ install_var_field (tree var, bool by_ref, int mask, omp_context *ctx,
     }
 
   if (mask & 1)
-    splay_tree_insert (ctx->field_map, key, (splay_tree_value) field);
+    {
+      splay_tree_insert (ctx->field_map, key, (splay_tree_value) field);
+      install_parm_decl (var, type, ctx);
+    }
   if ((mask & 2) && ctx->sfield_map)
     splay_tree_insert (ctx->sfield_map, key, (splay_tree_value) sfield);
 }
@@ -1088,6 +1132,8 @@ delete_omp_context (splay_tree_value value)
     splay_tree_delete (ctx->field_map);
   if (ctx->sfield_map)
     splay_tree_delete (ctx->sfield_map);
+  if (ctx->parm_map)
+    splay_tree_delete (ctx->parm_map);
 
   /* We hijacked DECL_ABSTRACT_ORIGIN earlier.  We need to clear it before
      it produces corrupt debug information.  */
@@ -1571,6 +1617,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx,
 		  insert_field_into_struct (ctx->record_type, field);
 		  splay_tree_insert (ctx->field_map, (splay_tree_key) decl,
 				     (splay_tree_value) field);
+		  install_parm_decl (decl, ptr_type_node, ctx);
 		}
 	    }
 	  break;
@@ -1849,10 +1896,13 @@ omp_maybe_offloaded_ctx (omp_context *ctx)
 }
 
 /* Build a decl for the omp child function.  It'll not contain a body
-   yet, just the bare decl.  */
+   yet, just the bare decl.  Unlike omp child functions, acc child
+   functions for parallel regions have one argument per data
+   mapping.  */
 
 static void
-create_omp_child_function (omp_context *ctx, bool task_copy)
+create_omp_child_function (omp_context *ctx, bool task_copy,
+			   unsigned int map_cnt = 0)
 {
   tree decl, type, name, t;
 
@@ -1860,6 +1910,13 @@ create_omp_child_function (omp_context *ctx, bool task_copy)
   if (task_copy)
     type = build_function_type_list (void_type_node, ptr_type_node,
 				     ptr_type_node, NULL_TREE);
+  else if (is_oacc_parallel (ctx))
+    {
+      tree *arg_types = (tree *) alloca (sizeof (tree) * map_cnt);
+      for (unsigned int i = 0; i < map_cnt; i++)
+	arg_types[i] = ptr_type_node;
+      type = build_function_type_array (void_type_node, map_cnt, arg_types);
+    }
   else
     type = build_function_type_list (void_type_node, ptr_type_node, NULL_TREE);
 
@@ -1933,33 +1990,35 @@ create_omp_child_function (omp_context *ctx, bool task_copy)
   DECL_CONTEXT (t) = decl;
   DECL_RESULT (decl) = t;
 
-  tree data_name = get_identifier (".omp_data_i");
-  t = build_decl (DECL_SOURCE_LOCATION (decl), PARM_DECL, data_name,
-		  ptr_type_node);
-  DECL_ARTIFICIAL (t) = 1;
-  DECL_NAMELESS (t) = 1;
-  DECL_ARG_TYPE (t) = ptr_type_node;
-  DECL_CONTEXT (t) = current_function_decl;
-  TREE_USED (t) = 1;
-  TREE_READONLY (t) = 1;
-  DECL_ARGUMENTS (decl) = t;
-  if (!task_copy)
-    ctx->receiver_decl = t;
-  else
+  if (!is_oacc_parallel (ctx))
     {
-      t = build_decl (DECL_SOURCE_LOCATION (decl),
-		      PARM_DECL, get_identifier (".omp_data_o"),
+      tree data_name = get_identifier (".omp_data_i");
+      t = build_decl (DECL_SOURCE_LOCATION (decl), PARM_DECL, data_name,
 		      ptr_type_node);
       DECL_ARTIFICIAL (t) = 1;
       DECL_NAMELESS (t) = 1;
       DECL_ARG_TYPE (t) = ptr_type_node;
       DECL_CONTEXT (t) = current_function_decl;
       TREE_USED (t) = 1;
-      TREE_ADDRESSABLE (t) = 1;
-      DECL_CHAIN (t) = DECL_ARGUMENTS (decl);
+      TREE_READONLY (t) = 1;
       DECL_ARGUMENTS (decl) = t;
+      if (!task_copy)
+	ctx->receiver_decl = t;
+      else
+	{
+	  t = build_decl (DECL_SOURCE_LOCATION (decl),
+			  PARM_DECL, get_identifier (".omp_data_o"),
+			  ptr_type_node);
+	  DECL_ARTIFICIAL (t) = 1;
+	  DECL_NAMELESS (t) = 1;
+	  DECL_ARG_TYPE (t) = ptr_type_node;
+	  DECL_CONTEXT (t) = current_function_decl;
+	  TREE_USED (t) = 1;
+	  TREE_ADDRESSABLE (t) = 1;
+	  DECL_CHAIN (t) = DECL_ARGUMENTS (decl);
+	  DECL_ARGUMENTS (decl) = t;
+	}
     }
-
   /* Allocate memory for the function structure.  The call to
      allocate_struct_function clobbers CFUN, so we need to restore
      it afterward.  */
@@ -2852,6 +2911,7 @@ scan_omp_target (gomp_target *stmt, omp_context *outer_ctx)
 
   ctx = new_omp_context (stmt, outer_ctx);
   ctx->field_map = splay_tree_new (splay_tree_compare_pointers, 0, 0);
+  ctx->parm_map = splay_tree_new (splay_tree_compare_pointers, 0, 0);
   ctx->record_type = lang_hooks.types.make_type (RECORD_TYPE);
   name = create_tmp_var_name (".omp_data_t");
   name = build_decl (gimple_location (stmt),
@@ -2864,8 +2924,11 @@ scan_omp_target (gomp_target *stmt, omp_context *outer_ctx)
   bool base_pointers_restrict = false;
   if (offloaded)
     {
-      create_omp_child_function (ctx, false);
-      gimple_omp_target_set_child_fn (stmt, ctx->cb.dst_fn);
+      if (!is_oacc_parallel (ctx))
+	{
+	  create_omp_child_function (ctx, false);
+	  gimple_omp_target_set_child_fn (stmt, ctx->cb.dst_fn);
+	}
 
       base_pointers_restrict = omp_target_base_pointers_restrict_p (clauses);
       if (base_pointers_restrict && dump_file && (dump_flags & TDF_DETAILS))
@@ -9693,6 +9756,18 @@ convert_from_firstprivate_int (tree var, tree orig_type, bool is_ref,
   return fold_build1 (VIEW_CONVERT_EXPR, type, tmp);
 }
 
+static tree
+append_decl_arg (tree var, tree decl_args, omp_context *ctx)
+{
+  if (!is_oacc_parallel (ctx))
+    return NULL_TREE;
+
+  tree temp = lookup_parm (var, ctx);
+  DECL_CHAIN (temp) = decl_args;
+
+  return temp;
+}
+
 /* Lower the GIMPLE_OMP_TARGET in the current statement
    in GSI_P.  CTX holds context information for the directive.  */
 
@@ -9706,7 +9781,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
   gimple_seq tgt_body, olist, ilist, fplist, new_body;
   location_t loc = gimple_location (stmt);
   bool offloaded, data_region;
-  unsigned int map_cnt = 0;
+  unsigned int map_cnt = 0, init_cnt = 0;
 
   offloaded = is_gimple_omp_offloaded (stmt);
   switch (gimple_omp_target_kind (stmt))
@@ -9754,11 +9829,83 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
     }
   else if (data_region)
     tgt_body = gimple_omp_body (stmt);
-  child_fn = ctx->cb.dst_fn;
 
   push_gimplify_context ();
   fplist = NULL;
 
+  /* Determine init_cnt to finish initialize ctx.  */
+
+  if (is_oacc_parallel (ctx))
+    {
+      for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c))
+	switch (OMP_CLAUSE_CODE (c))
+	  {
+	    tree var;
+
+	  default:
+	    break;
+	  case OMP_CLAUSE_MAP:
+	  case OMP_CLAUSE_TO:
+	  case OMP_CLAUSE_FROM:
+	  init_oacc_firstprivate:
+	    var = OMP_CLAUSE_DECL (c);
+	    if (!DECL_P (var))
+	      {
+		if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
+		    || (!OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c)
+			&& (OMP_CLAUSE_MAP_KIND (c)
+			    != GOMP_MAP_FIRSTPRIVATE_POINTER)))
+		  init_cnt++;
+		continue;
+	      }
+
+	    if (DECL_SIZE (var)
+		&& TREE_CODE (DECL_SIZE (var)) != INTEGER_CST)
+	      {
+		tree var2 = DECL_VALUE_EXPR (var);
+		gcc_assert (TREE_CODE (var2) == INDIRECT_REF);
+		var2 = TREE_OPERAND (var2, 0);
+		gcc_assert (DECL_P (var2));
+		var = var2;
+	      }
+
+	    if (offloaded
+		&& OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+		&& (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
+		    || (OMP_CLAUSE_MAP_KIND (c)
+			== GOMP_MAP_FIRSTPRIVATE_REFERENCE)))
+	      {
+		continue;
+	      }
+
+	    if (!maybe_lookup_field (var, ctx))
+	      continue;
+
+	    init_cnt++;
+	    break;
+
+	  case OMP_CLAUSE_FIRSTPRIVATE:
+	    if (is_oacc_parallel (ctx))
+	      goto init_oacc_firstprivate;
+	    init_cnt++;
+	    break;
+
+	  case OMP_CLAUSE_USE_DEVICE_PTR:
+	  case OMP_CLAUSE_IS_DEVICE_PTR:
+	    init_cnt++;
+	    break;
+	  }
+
+      /* Initialize the offloaded child function.  */
+
+      create_omp_child_function (ctx, false, init_cnt);
+      gimple_omp_target_set_child_fn (stmt, ctx->cb.dst_fn);
+    }
+
+  child_fn = ctx->cb.dst_fn;
+
+  /* Clause Pass 1: Scan and prepare sender decls VALUE_EXPRs for
+     usage on the child function.  */
   for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c))
     switch (OMP_CLAUSE_CODE (c))
       {
@@ -10019,6 +10166,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 
   if (offloaded)
     {
+      if (is_oacc_parallel (ctx))
+	gcc_assert (init_cnt == map_cnt);
       target_nesting_level++;
       lower_omp (&tgt_body, ctx);
       target_nesting_level--;
@@ -10068,6 +10217,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
       vec_alloc (vsize, map_cnt);
       vec_alloc (vkind, map_cnt);
       unsigned int map_idx = 0;
+      tree decl_args = NULL_TREE;
 
       for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c))
 	switch (OMP_CLAUSE_CODE (c))
@@ -10255,6 +10405,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 	    if (s == NULL_TREE)
 	      s = TYPE_SIZE_UNIT (TREE_TYPE (ovar));
 	    s = fold_convert (size_type_node, s);
+	    decl_args = append_decl_arg (ovar, decl_args, ctx);
 	    purpose = size_int (map_idx++);
 	    CONSTRUCTOR_APPEND_ELT (vsize, purpose, s);
 	    if (TREE_CODE (s) != INTEGER_CST)
@@ -10395,6 +10546,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 	    else
 	      s = TYPE_SIZE_UNIT (TREE_TYPE (ovar));
 	    s = fold_convert (size_type_node, s);
+	    decl_args = append_decl_arg (ovar, decl_args, ctx);
 	    purpose = size_int (map_idx++);
 	    CONSTRUCTOR_APPEND_ELT (vsize, purpose, s);
 	    if (TREE_CODE (s) != INTEGER_CST)
@@ -10467,6 +10619,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 				       gimple_build_label (opt_arg_label));
 	      }
 	    s = size_int (0);
+	    decl_args = append_decl_arg (ovar, decl_args, ctx);
 	    purpose = size_int (map_idx++);
 	    CONSTRUCTOR_APPEND_ELT (vsize, purpose, s);
 	    gcc_checking_assert (tkind
@@ -10479,6 +10632,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 	  }
 
       gcc_assert (map_idx == map_cnt);
+      if (is_oacc_parallel (ctx))
+	DECL_ARGUMENTS (child_fn) = nreverse (decl_args);
 
       DECL_INITIAL (TREE_VEC_ELT (t, 1))
 	= build_constructor (TREE_TYPE (TREE_VEC_ELT (t, 1)), vsize);
@@ -10517,9 +10672,12 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
     {
       t = build_fold_addr_expr_loc (loc, ctx->sender_decl);
       /* fixup_child_record_type might have changed receiver_decl's type.  */
-      t = fold_convert_loc (loc, TREE_TYPE (ctx->receiver_decl), t);
-      gimple_seq_add_stmt (&new_body,
-	  		   gimple_build_assign (ctx->receiver_decl, t));
+      if (!is_oacc_parallel (ctx))
+	{
+	  t = fold_convert_loc (loc, TREE_TYPE (ctx->receiver_decl), t);
+	  gimple_seq_add_stmt (&new_body,
+			       gimple_build_assign (ctx->receiver_decl, t));
+	}
     }
   gimple_seq_add_seq (&new_body, fplist);
 
diff --git a/gcc/testsuite/ChangeLog.omp b/gcc/testsuite/ChangeLog.omp
index c9917439f699..9bc3aac58f6d 100644
--- a/gcc/testsuite/ChangeLog.omp
+++ b/gcc/testsuite/ChangeLog.omp
@@ -1,3 +1,7 @@
+2017-12-21  Cesar Philippidis  <cesar@codesourcery.com>
+
+	* c-c++-common/goacc/large_array.c: New test.
+
 2018-12-21  Gergö Barany  <gergo@codesourcery.com>
 
 	* c-c++-common/goacc/host_data-1.c: Add tests of if and if_present
diff --git a/gcc/testsuite/c-c++-common/goacc/large_array.c b/gcc/testsuite/c-c++-common/goacc/large_array.c
new file mode 100644
index 000000000000..ce0f4c12a749
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/large_array.c
@@ -0,0 +1,18 @@
+/* Ensure that alloca'ed arrays can be transferred to the
+   accelerator. */
+
+/* { dg-require-effective-target alloca } */
+
+int
+main ()
+{
+  int n = 100, m = 10, i, j;
+  float a[n][m];
+
+  #pragma acc parallel loop
+  for (i = 0; i < n; i++)
+    for (j = 0; j < m; j++)
+      a[i][j] = 0;
+
+  return 0;
+}
diff --git a/gcc/tree-ssa-structalias.c b/gcc/tree-ssa-structalias.c
index 4a0b02e9b03c..2f29c39565e4 100644
--- a/gcc/tree-ssa-structalias.c
+++ b/gcc/tree-ssa-structalias.c
@@ -4692,6 +4692,7 @@ find_func_aliases_for_builtin_call (struct function *fn, gcall *t)
       case BUILT_IN_GOMP_PARALLEL:
       case BUILT_IN_GOACC_PARALLEL:
 	{
+	  bool oacc_parallel = false;
 	  if (in_ipa_mode)
 	    {
 	      unsigned int fnpos, argpos;
@@ -4705,13 +4706,17 @@ find_func_aliases_for_builtin_call (struct function *fn, gcall *t)
 		case BUILT_IN_GOACC_PARALLEL:
 		  /* __builtin_GOACC_parallel (flags_m, fn, mapnum, hostaddrs,
 					       sizes, kinds, ...).  */
-		  fnpos = 1;
-		  argpos = 3;
+		  fnpos = 2;
+		  argpos = 4;
+		  oacc_parallel = gimple_call_arg (t, 1) == integer_one_node;
 		  break;
 		default:
 		  gcc_unreachable ();
 		}
 
+	      if (oacc_parallel)
+		break;
+
 	      tree fnarg = gimple_call_arg (t, fnpos);
 	      gcc_assert (TREE_CODE (fnarg) == ADDR_EXPR);
 	      tree fndecl = TREE_OPERAND (fnarg, 0);
@@ -5253,6 +5258,7 @@ find_func_clobbers (struct function *fn, gimple *origt)
 	      unsigned int fnpos, argpos;
 	      unsigned int implicit_use_args[2];
 	      unsigned int num_implicit_use_args = 0;
+	      bool oacc_parallel = false;
 	      switch (DECL_FUNCTION_CODE (decl))
 		{
 		case BUILT_IN_GOMP_PARALLEL:
@@ -5263,15 +5269,19 @@ find_func_clobbers (struct function *fn, gimple *origt)
 		case BUILT_IN_GOACC_PARALLEL:
 		  /* __builtin_GOACC_parallel (flags_m, fn, mapnum, hostaddrs,
 					       sizes, kinds, ...).  */
-		  fnpos = 1;
-		  argpos = 3;
-		  implicit_use_args[num_implicit_use_args++] = 4;
+		  fnpos = 2;
+		  argpos = 4;
 		  implicit_use_args[num_implicit_use_args++] = 5;
+		  implicit_use_args[num_implicit_use_args++] = 6;
+		  oacc_parallel = gimple_call_arg (t, 1) == integer_one_node;
 		  break;
 		default:
 		  gcc_unreachable ();
 		}
 
+	      if (oacc_parallel)
+		break;
+
 	      tree fnarg = gimple_call_arg (t, fnpos);
 	      gcc_assert (TREE_CODE (fnarg) == ADDR_EXPR);
 	      tree fndecl = TREE_OPERAND (fnarg, 0);
@@ -8206,7 +8216,7 @@ ipa_pta_execute (void)
 		if (gimple_call_builtin_p (stmt, BUILT_IN_GOMP_PARALLEL))
 		  called_decl = TREE_OPERAND (gimple_call_arg (stmt, 0), 0);
 		else if (gimple_call_builtin_p (stmt, BUILT_IN_GOACC_PARALLEL))
-		  called_decl = TREE_OPERAND (gimple_call_arg (stmt, 1), 0);
+		  called_decl = TREE_OPERAND (gimple_call_arg (stmt, 2), 0);
 
 		if (called_decl != NULL_TREE
 		    && !fndecl_maybe_in_other_partition (called_decl))
diff --git a/libgomp/ChangeLog.omp b/libgomp/ChangeLog.omp
index 5dc06978255f..05cfabc72af6 100644
--- a/libgomp/ChangeLog.omp
+++ b/libgomp/ChangeLog.omp
@@ -1,3 +1,43 @@
+2017-12-21  Cesar Philippidis  <cesar@codesourcery.com>
+
+	* Makefile.am: Add libffi build dependency.
+	* configure.ac: Likewise.
+	* Makefile.in: Regenerate.
+	* config.h.in: Regenerate.
+	* configure: Regenerate.
+	* libgomp-plugin.h: Define GOMP_OFFLOAD_openacc_exec_params and
+	GOMP_OFFLOAD_openacc_async_exec_params.
+	* libgomp.h (acc_dispatch_t): Use them here.
+	* libgomp.map (GOACC_parallel_keyed_v2): Declare.
+	* libgomp_g.h (GOACC_parallel_keyed_v2): Likewise.
+	* oacc-host.c (host_openacc_exec_params): New function.
+	(host_openacc_async_exec_params): Likewise.
+	* oacc-parallel.c (goacc_call_host_fn): Likewise.
+	(GOACC_parallel_keyed_internal): Likewise.
+	(GOACC_parallel_keyed): Wrapper for GOACC_parallel_keyed_internal.
+	(GOACC_parallel_keyed_v2): Likewise.
+	* plugin/plugin-nvptx.c (nvptx_exec): Replace CUDeviceptr dp parameter
+	with void **kargs.
+	(openacc_exec_internal): New function.
+	(GOMP_OFFLOAD_openacc_exec_params): New function.
+	(GOMP_OFFLOAD_openacc_exec): Update to call openacc_exec_internal.
+	(openacc_async_exec_internal): New function.
+	(GOMP_OFFLOAD_openacc_async_exec_params): New function.
+	(GOMP_OFFLOAD_openacc_async_exec): Update call to
+	openacc_async_exec_internal.
+	* target.c (gomp_load_plugin_for_device): Handle
+	openacc_exec_params and openacc_async_exec_params.
+	* testsuite/Makefile.in: Regenerate.
+	* testsuite/libgomp.oacc-c-c++-common/combined-directives-1.c:
+	Xfail on offloaded targets.
+
+	* Makefile.def: Bootstrap module libffi. Add libffi dependency
+	to all-target-libgomp.
+	* Makefile.in: Regenerate.
+	* configure.ac: Add libffi to bootstrap_target_libs when libgomp
+	is bootstrapped.
+	* configure: Regenerate.
+
 2018-12-21  Gergö Barany  <gergo@codesourcery.com>
 
 	* libgomp.h (enum gomp_map_vars_kind): Add
diff --git a/libgomp/Makefile.am b/libgomp/Makefile.am
index 062fded0805a..2db386531c3e 100644
--- a/libgomp/Makefile.am
+++ b/libgomp/Makefile.am
@@ -14,9 +14,16 @@ search_path = $(addprefix $(top_srcdir)/config/, $(config_path)) $(top_srcdir) \
 fincludedir = $(libdir)/gcc/$(target_alias)/$(gcc_version)$(MULTISUBDIR)/finclude
 libsubincludedir = $(libdir)/gcc/$(target_alias)/$(gcc_version)/include
 
+LIBFFI = @LIBFFI@
+LIBFFIINCS = @LIBFFIINCS@
+
+if USE_LIBFFI
+libgomp_la_LIBADD = $(LIBFFI)
+endif
+
 vpath % $(strip $(search_path))
 
-AM_CPPFLAGS = $(addprefix -I, $(search_path))
+AM_CPPFLAGS = $(addprefix -I, $(search_path)) $(LIBFFIINCS)
 AM_CFLAGS = $(XCFLAGS)
 AM_LDFLAGS = $(XLDFLAGS) $(SECTION_LDFLAGS) $(OPT_LDFLAGS)
 
diff --git a/libgomp/Makefile.in b/libgomp/Makefile.in
index c7e63f120fdb..820a576e9ed2 100644
--- a/libgomp/Makefile.in
+++ b/libgomp/Makefile.in
@@ -207,7 +207,6 @@ libgomp_plugin_nvptx_la_LINK = $(LIBTOOL) $(AM_V_lt) --tag=CC \
 	$(libgomp_plugin_nvptx_la_LDFLAGS) $(LDFLAGS) -o $@
 @PLUGIN_NVPTX_TRUE@am_libgomp_plugin_nvptx_la_rpath = -rpath \
 @PLUGIN_NVPTX_TRUE@	$(toolexeclibdir)
-libgomp_la_LIBADD =
 @USE_FORTRAN_TRUE@am__objects_1 = openacc.lo
 am_libgomp_la_OBJECTS = alloc.lo atomic.lo barrier.lo critical.lo \
 	env.lo error.lo icv.lo icv-device.lo iter.lo iter_ull.lo \
@@ -390,6 +389,8 @@ INSTALL_SCRIPT = @INSTALL_SCRIPT@
 INSTALL_STRIP_PROGRAM = @INSTALL_STRIP_PROGRAM@
 LD = @LD@
 LDFLAGS = @LDFLAGS@
+LIBFFI = @LIBFFI@
+LIBFFIINCS = @LIBFFIINCS@
 LIBOBJS = @LIBOBJS@
 LIBS = @LIBS@
 LIBTOOL = @LIBTOOL@
@@ -523,7 +524,8 @@ search_path = $(addprefix $(top_srcdir)/config/, $(config_path)) $(top_srcdir) \
 
 fincludedir = $(libdir)/gcc/$(target_alias)/$(gcc_version)$(MULTISUBDIR)/finclude
 libsubincludedir = $(libdir)/gcc/$(target_alias)/$(gcc_version)/include
-AM_CPPFLAGS = $(addprefix -I, $(search_path))
+libgomp_la_LIBADD = $(LIBFFI)
+AM_CPPFLAGS = $(addprefix -I, $(search_path)) $(LIBFFIINCS)
 AM_CFLAGS = $(XCFLAGS)
 AM_LDFLAGS = $(XLDFLAGS) $(SECTION_LDFLAGS) $(OPT_LDFLAGS)
 toolexeclib_LTLIBRARIES = libgomp.la $(am__append_1) $(am__append_2)
diff --git a/libgomp/config.h.in b/libgomp/config.h.in
index 449cc8506878..5ad5f532af19 100644
--- a/libgomp/config.h.in
+++ b/libgomp/config.h.in
@@ -210,5 +210,8 @@
 /* Define to 1 if the target use emutls for thread-local storage. */
 #undef USE_EMUTLS
 
+/* Define to 1 if the target requires libffi to call the offloaded funtions. */
+#undef USE_LIBFFI
+
 /* Version number of package */
 #undef VERSION
diff --git a/libgomp/configure b/libgomp/configure
index 529a2b4711d5..032620bf8c5e 100755
--- a/libgomp/configure
+++ b/libgomp/configure
@@ -681,6 +681,10 @@ PLUGIN_NVPTX_CPPFLAGS
 PLUGIN_NVPTX
 CUDA_DRIVER_LIB
 CUDA_DRIVER_INCLUDE
+USE_LIBFFI_FALSE
+USE_LIBFFI_TRUE
+LIBFFIINCS
+LIBFFI
 libtool_VERSION
 ac_ct_FC
 FCFLAGS
@@ -2704,7 +2708,6 @@ else
 fi
 
 
-
 # -------
 # -------
 
@@ -11393,7 +11396,7 @@ else
   lt_dlunknown=0; lt_dlno_uscore=1; lt_dlneed_uscore=2
   lt_status=$lt_dlunknown
   cat > conftest.$ac_ext <<_LT_EOF
-#line 11396 "configure"
+#line 11399 "configure"
 #include "confdefs.h"
 
 #if HAVE_DLFCN_H
@@ -11499,7 +11502,7 @@ else
   lt_dlunknown=0; lt_dlno_uscore=1; lt_dlneed_uscore=2
   lt_status=$lt_dlunknown
   cat > conftest.$ac_ext <<_LT_EOF
-#line 11502 "configure"
+#line 11505 "configure"
 #include "confdefs.h"
 
 #if HAVE_DLFCN_H
@@ -15374,9 +15377,31 @@ $as_echo "#define LIBGOMP_OFFLOADED_ONLY 1" >>confdefs.h
 
 fi
 
+# Prepare libffi when necessary.
+
+LIBFFI=
+LIBFFIINCS=
+if test -d ../libffi; then
+
+$as_echo "#define USE_LIBFFI 1" >>confdefs.h
+
+   LIBFFI=../libffi/libffi_convenience.la
+   LIBFFIINCS='-I$(top_srcdir)/../libffi/include -I../libffi/include'
+fi
+
+
+ if test -d ../libffi; then
+  USE_LIBFFI_TRUE=
+  USE_LIBFFI_FALSE='#'
+else
+  USE_LIBFFI_TRUE='#'
+  USE_LIBFFI_FALSE=
+fi
+
+
 # Plugins for offload execution, configure.ac fragment.  -*- mode: autoconf -*-
 #
-# Copyright (C) 2014-2018 Free Software Foundation, Inc.
+# Copyright (C) 2014-2019 Free Software Foundation, Inc.
 #
 # Contributed by Mentor Embedded.
 #
@@ -15824,8 +15849,7 @@ for ac_func in aligned_alloc posix_memalign memalign _aligned_malloc
 do :
   as_ac_var=`$as_echo "ac_cv_func_$ac_func" | $as_tr_sh`
 ac_fn_c_check_func "$LINENO" "$ac_func" "$as_ac_var"
-eval as_val=\$$as_ac_var
-   if test "x$as_val" = x""yes; then :
+if eval test \"x\$"$as_ac_var"\" = x"yes"; then :
   cat >>confdefs.h <<_ACEOF
 #define `$as_echo "HAVE_$ac_func" | $as_tr_cpp` 1
 _ACEOF
@@ -17409,6 +17433,10 @@ if test -z "${MAINTAINER_MODE_TRUE}" && test -z "${MAINTAINER_MODE_FALSE}"; then
   as_fn_error $? "conditional \"MAINTAINER_MODE\" was never defined.
 Usually this means the macro was only invoked conditionally." "$LINENO" 5
 fi
+if test -z "${USE_LIBFFI_TRUE}" && test -z "${USE_LIBFFI_FALSE}"; then
+  as_fn_error $? "conditional \"USE_LIBFFI\" was never defined.
+Usually this means the macro was only invoked conditionally." "$LINENO" 5
+fi
 if test -z "${PLUGIN_NVPTX_TRUE}" && test -z "${PLUGIN_NVPTX_FALSE}"; then
   as_fn_error $? "conditional \"PLUGIN_NVPTX\" was never defined.
 Usually this means the macro was only invoked conditionally." "$LINENO" 5
diff --git a/libgomp/configure.ac b/libgomp/configure.ac
index f75c6226566c..2350b78f310e 100644
--- a/libgomp/configure.ac
+++ b/libgomp/configure.ac
@@ -27,7 +27,6 @@ LIBGOMP_ENABLE(generated-files-in-srcdir, no, ,
 AC_MSG_RESULT($enable_generated_files_in_srcdir)
 AM_CONDITIONAL(GENINSRC, test "$enable_generated_files_in_srcdir" = yes)
 
-
 # -------
 # -------
 
@@ -214,6 +213,19 @@ if test x$libgomp_offloaded_only = xyes; then
             [Define to 1 if building libgomp for an accelerator-only target.])
 fi
 
+# Prepare libffi when necessary.
+
+LIBFFI=
+LIBFFIINCS=
+if test -d ../libffi; then
+   AC_DEFINE(USE_LIBFFI, 1, [Define if we're to use libffi.])
+   LIBFFI=../libffi/libffi_convenience.la
+   LIBFFIINCS='-I$(top_srcdir)/../libffi/include -I../libffi/include'
+fi
+AC_SUBST(LIBFFI)
+AC_SUBST(LIBFFIINCS)
+AM_CONDITIONAL([USE_LIBFFI], [test -d ../libffi])
+
 m4_include([plugin/configfrag.ac])
 
 # Check for functions needed.
diff --git a/libgomp/libgomp-plugin.h b/libgomp/libgomp-plugin.h
index 01483f27f4cb..be7c994faebc 100644
--- a/libgomp/libgomp-plugin.h
+++ b/libgomp/libgomp-plugin.h
@@ -123,6 +123,13 @@ extern void GOMP_OFFLOAD_openacc_async_queue_callback (struct goacc_asyncqueue *
 extern void GOMP_OFFLOAD_openacc_async_exec (void (*) (void *), size_t, void **,
 					     void **, unsigned *, void *,
 					     struct goacc_asyncqueue *);
+extern void GOMP_OFFLOAD_openacc_exec_params (void (*) (void *), size_t,
+					      void **, void **, unsigned *,
+					      void *);
+extern void GOMP_OFFLOAD_openacc_async_exec_params (void (*) (void *), size_t,
+						    void **, void **,
+						    unsigned *, void *,
+						    struct goacc_asyncqueue *);
 extern bool GOMP_OFFLOAD_openacc_async_dev2host (int, void *, const void *, size_t,
 						 struct goacc_asyncqueue *);
 extern bool GOMP_OFFLOAD_openacc_async_host2dev (int, void *, const void *, size_t,
diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h
index 6bc39ce9aae4..9c9157d826c1 100644
--- a/libgomp/libgomp.h
+++ b/libgomp/libgomp.h
@@ -975,6 +975,7 @@ typedef struct acc_dispatch_t
 {
   /* Execute.  */
   __typeof (GOMP_OFFLOAD_openacc_exec) *exec_func;
+  __typeof (GOMP_OFFLOAD_openacc_exec_params) *exec_params_func;
 
   /* Create/destroy TLS data.  */
   __typeof (GOMP_OFFLOAD_openacc_create_thread_data) *create_thread_data_func;
@@ -998,6 +999,7 @@ typedef struct acc_dispatch_t
     __typeof (GOMP_OFFLOAD_openacc_async_queue_callback) *queue_callback_func;
 
     __typeof (GOMP_OFFLOAD_openacc_async_exec) *exec_func;
+    __typeof (GOMP_OFFLOAD_openacc_async_exec_params) *exec_params_func;
     __typeof (GOMP_OFFLOAD_openacc_async_dev2host) *dev2host_func;
     __typeof (GOMP_OFFLOAD_openacc_async_host2dev) *host2dev_func;
   } async;
diff --git a/libgomp/libgomp.map b/libgomp/libgomp.map
index f662bc864d65..aa76ee309ffd 100644
--- a/libgomp/libgomp.map
+++ b/libgomp/libgomp.map
@@ -507,8 +507,10 @@ GOACC_2.0.1 {
 GOACC_2.0.GOMP_4_BRANCH {
   global:
 	GOMP_set_offload_targets;
+	GOACC_parallel_keyed_v2;
 } GOACC_2.0.1;
 
+
 GOMP_PLUGIN_1.0 {
   global:
 	GOMP_PLUGIN_malloc;
diff --git a/libgomp/libgomp_g.h b/libgomp/libgomp_g.h
index 5c4bd8d34971..c3173417bd90 100644
--- a/libgomp/libgomp_g.h
+++ b/libgomp/libgomp_g.h
@@ -362,6 +362,8 @@ extern void GOMP_teams_reg (void (*) (void *), void *, unsigned, unsigned,
 
 extern void GOACC_parallel_keyed (int, void (*) (void *), size_t,
 				  void **, size_t *, unsigned short *, ...);
+extern void GOACC_parallel_keyed_v2 (int, int, void (*) (void *), size_t,
+				  void **, size_t *, unsigned short *, ...);
 extern void GOACC_parallel (int, void (*) (void *), size_t, void **, size_t *,
 			    unsigned short *, int, int, int, int, int, ...);
 extern void GOACC_data_start (int, size_t, void **, size_t *,
diff --git a/libgomp/oacc-host.c b/libgomp/oacc-host.c
index 78de88e84942..b19b7479afd6 100644
--- a/libgomp/oacc-host.c
+++ b/libgomp/oacc-host.c
@@ -158,6 +158,30 @@ host_openacc_async_exec (void (*fn) (void *),
   fn (hostaddrs);
 }
 
+static void
+host_openacc_exec_params (void (*fn) (void *),
+			  size_t mapnum __attribute__ ((unused)),
+			  void **hostaddrs,
+			  void **devaddrs __attribute__ ((unused)),
+			  unsigned *dims __attribute__ ((unused)),
+			  void *targ_mem_desc __attribute__ ((unused)))
+{
+  fn (hostaddrs);
+}
+
+static void
+host_openacc_async_exec_params (void (*fn) (void *),
+				size_t mapnum __attribute__ ((unused)),
+				void **hostaddrs,
+				void **devaddrs __attribute__ ((unused)),
+				unsigned *dims __attribute__ ((unused)),
+				void *targ_mem_desc __attribute__ ((unused)),
+				struct goacc_asyncqueue *aq __attribute__ ((unused)))
+{
+  fn (hostaddrs);
+}
+
+
 static int
 host_openacc_async_test (struct goacc_asyncqueue *aq __attribute__ ((unused)))
 {
@@ -266,6 +290,7 @@ static struct gomp_device_descr host_dispatch =
 
     .openacc = {
       .exec_func = host_openacc_exec,
+      .exec_params_func = host_openacc_exec_params,
 
       .create_thread_data_func = host_openacc_create_thread_data,
       .destroy_thread_data_func = host_openacc_destroy_thread_data,
@@ -278,6 +303,7 @@ static struct gomp_device_descr host_dispatch =
 	.serialize_func = host_openacc_async_serialize,
 	.queue_callback_func = host_openacc_async_queue_callback,
 	.exec_func = host_openacc_async_exec,
+	.exec_params_func = host_openacc_async_exec_params,
 	.dev2host_func = host_openacc_async_dev2host,
 	.host2dev_func = host_openacc_async_host2dev,
       },
diff --git a/libgomp/oacc-parallel.c b/libgomp/oacc-parallel.c
index 98302f1f7adb..b949599a8d00 100644
--- a/libgomp/oacc-parallel.c
+++ b/libgomp/oacc-parallel.c
@@ -31,6 +31,9 @@
 #include "libgomp_g.h"
 #include "gomp-constants.h"
 #include "oacc-int.h"
+#if USE_LIBFFI
+# include "ffi.h"
+#endif
 #ifdef HAVE_INTTYPES_H
 # include <inttypes.h>  /* For PRIu64.  */
 #endif
@@ -136,20 +139,48 @@ handle_ftn_pointers (size_t mapnum, void **hostaddrs, size_t *sizes,
 
 static void goacc_wait (int async, int num_waits, va_list *ap);
 
+static void
+goacc_call_host_fn (void (*fn) (void *), size_t mapnum, void **hostaddrs,
+		    int params)
+{
+#ifdef USE_LIBFFI
+  ffi_cif cif;
+  ffi_type *arg_types[mapnum];
+  void *arg_values[mapnum];
+  ffi_arg result;
+  int i;
+
+  if (params)
+    {
+      for (i = 0; i < mapnum; i++)
+	{
+	  arg_types[i] = &ffi_type_pointer;
+	  arg_values[i] = &hostaddrs[i];
+	}
+
+      if (ffi_prep_cif (&cif, FFI_DEFAULT_ABI, mapnum,
+			&ffi_type_void, arg_types) == FFI_OK)
+	ffi_call (&cif, FFI_FN (fn), &result, arg_values);
+      else
+	abort ();
+    }
+  else
+#endif
+  fn (hostaddrs);
+}
 
 /* Launch a possibly offloaded function with FLAGS.  FN is the host fn
    address.  MAPNUM, HOSTADDRS, SIZES & KINDS  describe the memory
    blocks to be copied to/from the device.  Varadic arguments are
    keyed optional parameters terminated with a zero.  */
 
-void
-GOACC_parallel_keyed (int flags_m, void (*fn) (void *),
-		      size_t mapnum, void **hostaddrs, size_t *sizes,
-		      unsigned short *kinds, ...)
+static void
+GOACC_parallel_keyed_internal (int flags_m, int params, void (*fn) (void *),
+			       size_t mapnum, void **hostaddrs, size_t *sizes,
+			       unsigned short *kinds, va_list *ap)
 {
   int flags = GOACC_FLAGS_UNMARSHAL (flags_m);
 
-  va_list ap;
   struct goacc_thread *thr;
   struct gomp_device_descr *acc_dev;
   struct target_mem_desc *tgt;
@@ -181,13 +212,13 @@ GOACC_parallel_keyed (int flags_m, void (*fn) (void *),
   if (flags & GOACC_FLAG_HOST_FALLBACK)
     {
       goacc_save_and_set_bind (acc_device_host);
-      fn (hostaddrs);
+      goacc_call_host_fn (fn, mapnum, hostaddrs, params);
       goacc_restore_bind ();
       return;
     }
   else if (acc_device_type (acc_dev->type) == acc_device_host)
     {
-      fn (hostaddrs);
+      goacc_call_host_fn (fn, mapnum, hostaddrs, params);
       return;
     }
 
@@ -195,9 +226,8 @@ GOACC_parallel_keyed (int flags_m, void (*fn) (void *),
   for (i = 0; i != GOMP_DIM_MAX; i++)
     dims[i] = 0;
 
-  va_start (ap, kinds);
   /* TODO: This will need amending when device_type is implemented.  */
-  while ((tag = va_arg (ap, unsigned)) != 0)
+  while ((tag = va_arg (*ap, unsigned)) != 0)
     {
       if (GOMP_LAUNCH_DEVICE (tag))
 	gomp_fatal ("device_type '%d' offload parameters, libgomp is too old",
@@ -211,7 +241,7 @@ GOACC_parallel_keyed (int flags_m, void (*fn) (void *),
 
 	    for (i = 0; i != GOMP_DIM_MAX; i++)
 	      if (mask & GOMP_DIM_MASK (i))
-		dims[i] = va_arg (ap, unsigned);
+		dims[i] = va_arg (*ap, unsigned);
 	  }
 	  break;
 
@@ -221,14 +251,14 @@ GOACC_parallel_keyed (int flags_m, void (*fn) (void *),
 	    async = GOMP_LAUNCH_OP (tag);
 
 	    if (async == GOMP_LAUNCH_OP_MAX)
-	      async = va_arg (ap, unsigned);
+	      async = va_arg (*ap, unsigned);
 	    break;
 	  }
 
 	case GOMP_LAUNCH_WAIT:
 	  {
 	    unsigned num_waits = GOMP_LAUNCH_OP (tag);
-	    goacc_wait (async, num_waits, &ap);
+	    goacc_wait (async, num_waits, ap);
 	    break;
 	  }
 
@@ -237,7 +267,6 @@ GOACC_parallel_keyed (int flags_m, void (*fn) (void *),
 		      " libgomp is too old", GOMP_LAUNCH_CODE (tag));
 	}
     }
-  va_end (ap);
   
   if (!(acc_dev->capabilities & GOMP_OFFLOAD_CAP_NATIVE_EXEC))
     {
@@ -275,15 +304,23 @@ GOACC_parallel_keyed (int flags_m, void (*fn) (void *),
 
   if (aq == NULL)
     {
-      acc_dev->openacc.exec_func (tgt_fn, mapnum, hostaddrs, devaddrs,
-				  dims, tgt);
+      if (params)
+	acc_dev->openacc.exec_params_func (tgt_fn, mapnum, hostaddrs, devaddrs,
+					   dims, tgt);
+      else
+	acc_dev->openacc.exec_func (tgt_fn, mapnum, hostaddrs, devaddrs,
+				    dims, tgt);
       /* If running synchronously, unmap immediately.  */
       gomp_unmap_vars (tgt, true);
     }
   else
     {
-      acc_dev->openacc.async.exec_func (tgt_fn, mapnum, hostaddrs, devaddrs,
-					dims, tgt, aq);
+      if (params)
+	acc_dev->openacc.async.exec_params_func (tgt_fn, mapnum, hostaddrs,
+						 devaddrs, dims, tgt, aq);
+      else
+	acc_dev->openacc.async.exec_func (tgt_fn, mapnum, hostaddrs,
+					  devaddrs, dims, tgt, aq);
       gomp_unmap_vars_async (tgt, true, aq);
     }
 
@@ -294,6 +331,30 @@ GOACC_parallel_keyed (int flags_m, void (*fn) (void *),
 #endif
 }
 
+void
+GOACC_parallel_keyed (int flags_m, void (*fn) (void *),
+		      size_t mapnum, void **hostaddrs, size_t *sizes,
+		      unsigned short *kinds, ...)
+{
+  va_list ap;
+  va_start (ap, kinds);
+  GOACC_parallel_keyed_internal (flags_m, 0, fn, mapnum, hostaddrs, sizes,
+				 kinds, &ap);
+  va_end (ap);
+}
+
+void
+GOACC_parallel_keyed_v2 (int flags_m, int args, void (*fn) (void *),
+			 size_t mapnum, void **hostaddrs, size_t *sizes,
+			 unsigned short *kinds, ...)
+{
+  va_list ap;
+  va_start (ap, kinds);
+  GOACC_parallel_keyed_internal (flags_m, args, fn, mapnum, hostaddrs, sizes,
+				 kinds, &ap);
+  va_end (ap);
+}
+
 /* Legacy entry point, only provide host execution.  */
 
 void
diff --git a/libgomp/plugin/plugin-nvptx.c b/libgomp/plugin/plugin-nvptx.c
index 8f71e69acb60..b7a1a6c40f57 100644
--- a/libgomp/plugin/plugin-nvptx.c
+++ b/libgomp/plugin/plugin-nvptx.c
@@ -696,12 +696,11 @@ link_ptx (CUmodule *module, const struct targ_ptx_obj *ptx_objs,
 static void
 nvptx_exec (void (*fn), size_t mapnum, void **hostaddrs, void **devaddrs,
 	    unsigned *dims, void *targ_mem_desc,
-	    CUdeviceptr dp, CUstream stream)
+	    void **kargs, CUstream stream)
 {
   struct targ_fn_descriptor *targ_fn = (struct targ_fn_descriptor *) fn;
   CUfunction function;
   int i;
-  void *kargs[1];
   struct nvptx_thread *nvthd = nvptx_thread ();
   int warp_size = nvthd->ptx_dev->warp_size;
 
@@ -904,7 +903,6 @@ nvptx_exec (void (*fn), size_t mapnum, void **hostaddrs, void **devaddrs,
   // num_gangs		nctaid.x
   // num_workers	ntid.y
   // vector length	ntid.x
-  kargs[0] = &dp;
   CUDA_CALL_ASSERT (cuLaunchKernel, function,
 		    dims[GOMP_DIM_GANG], 1, 1,
 		    dims[GOMP_DIM_VECTOR], dims[GOMP_DIM_WORKER], 1,
@@ -1243,32 +1241,46 @@ GOMP_OFFLOAD_free (int ord, void *ptr)
 	  && nvptx_free (ptr, ptx_devices[ord]));
 }
 
-void
-GOMP_OFFLOAD_openacc_exec (void (*fn) (void *), size_t mapnum,
-			   void **hostaddrs, void **devaddrs,
-			   unsigned *dims, void *targ_mem_desc)
+static void
+openacc_exec_internal (void (*fn) (void *), int params, size_t mapnum,
+		       void **hostaddrs, void **devaddrs,
+		       unsigned *dims, void *targ_mem_desc)
 {
   GOMP_PLUGIN_debug (0, "  %s: prepare mappings\n", __FUNCTION__);
 
-  void **hp = NULL;
+  void **hp = alloca (mapnum * sizeof (void *));
   CUdeviceptr dp = 0;
 
   if (mapnum > 0)
     {
-      hp = alloca (mapnum * sizeof (void *));
-      for (int i = 0; i < mapnum; i++)
-	hp[i] = (devaddrs[i] ? devaddrs[i] : hostaddrs[i]);
-      CUDA_CALL_ASSERT (cuMemAlloc, &dp, mapnum * sizeof (void *));
+      if (params)
+	{
+	  for (int i = 0; i < mapnum; i++)
+	    hp[i] = (devaddrs[i] ? &devaddrs[i] : &hostaddrs[i]);
+	}
+      else
+	{
+	  for (int i = 0; i < mapnum; i++)
+	    hp[i] = (devaddrs[i] ? devaddrs[i] : hostaddrs[i]);
+	  CUDA_CALL_ASSERT (cuMemAlloc, &dp, mapnum * sizeof (void *));
+	}
     }
 
   /* Copy the (device) pointers to arguments to the device (dp and hp might in
      fact have the same value on a unified-memory system).  */
-  if (mapnum > 0)
+  if (!params && mapnum > 0)
     CUDA_CALL_ASSERT (cuMemcpyHtoD, dp, (void *) hp,
 		      mapnum * sizeof (void *));
 
-  nvptx_exec (fn, mapnum, hostaddrs, devaddrs, dims, targ_mem_desc,
-	      dp, NULL);
+  if (params)
+    nvptx_exec (fn, mapnum, hostaddrs, devaddrs, dims, targ_mem_desc,
+		hp, NULL);
+  else
+    {
+      void *kargs[1] = { &dp };
+      nvptx_exec (fn, mapnum, hostaddrs, devaddrs, dims, targ_mem_desc,
+		  kargs, NULL);
+    }
 
   CUresult r = CUDA_CALL_NOCHECK (cuStreamSynchronize, NULL);
   const char *maybe_abort_msg = "(perhaps abort was called)";
@@ -1277,7 +1289,27 @@ GOMP_OFFLOAD_openacc_exec (void (*fn) (void *), size_t mapnum,
 		       maybe_abort_msg);
   else if (r != CUDA_SUCCESS)
     GOMP_PLUGIN_fatal ("cuStreamSynchronize error: %s", cuda_error (r));
-  CUDA_CALL_ASSERT (cuMemFree, dp);
+
+  if (!params)
+    CUDA_CALL_ASSERT (cuMemFree, dp);
+}
+
+void
+GOMP_OFFLOAD_openacc_exec_params (void (*fn) (void *), size_t mapnum,
+			   void **hostaddrs, void **devaddrs,
+			   unsigned *dims, void *targ_mem_desc)
+{
+  openacc_exec_internal (fn, 1, mapnum, hostaddrs, devaddrs, dims,
+			 targ_mem_desc);
+}
+
+void
+GOMP_OFFLOAD_openacc_exec (void (*fn) (void *), size_t mapnum,
+			   void **hostaddrs, void **devaddrs,
+			   unsigned *dims, void *targ_mem_desc)
+{
+  openacc_exec_internal (fn, 0, mapnum, hostaddrs, devaddrs, dims,
+			 targ_mem_desc);
 }
 
 static void
@@ -1288,11 +1320,11 @@ cuda_free_argmem (void *ptr)
   free (block);
 }
 
-void
-GOMP_OFFLOAD_openacc_async_exec (void (*fn) (void *), size_t mapnum,
-				 void **hostaddrs, void **devaddrs,
-				 unsigned *dims, void *targ_mem_desc,
-				 struct goacc_asyncqueue *aq)
+static void
+openacc_async_exec_internal (void (*fn) (void *), int params, size_t mapnum,
+			     void **hostaddrs, void **devaddrs,
+			     unsigned *dims, void *targ_mem_desc,
+			     struct goacc_asyncqueue *aq)
 {
   GOMP_PLUGIN_debug (0, "  %s: prepare mappings\n", __FUNCTION__);
 
@@ -1302,16 +1334,25 @@ GOMP_OFFLOAD_openacc_async_exec (void (*fn) (void *), size_t mapnum,
 
   if (mapnum > 0)
     {
-      block = (void **) GOMP_PLUGIN_malloc ((mapnum + 2) * sizeof (void *));
-      hp = block + 2;
-      for (int i = 0; i < mapnum; i++)
-	hp[i] = (devaddrs[i] ? devaddrs[i] : hostaddrs[i]);
-      CUDA_CALL_ASSERT (cuMemAlloc, &dp, mapnum * sizeof (void *));
+      if (params)
+	{
+	  hp = alloca (sizeof (void *) * mapnum);
+	  for (int i = 0; i < mapnum; i++)
+	    hp[i] = (devaddrs[i] ? &devaddrs[i] : &hostaddrs[i]);
+	}
+      else
+	{
+	  block = (void **) GOMP_PLUGIN_malloc ((mapnum + 2) * sizeof (void *));
+	  hp = block + 2;
+	  for (int i = 0; i < mapnum; i++)
+	    hp[i] = (devaddrs[i] ? devaddrs[i] : hostaddrs[i]);
+	  CUDA_CALL_ASSERT (cuMemAlloc, &dp, mapnum * sizeof (void *));
+	}
     }
 
   /* Copy the (device) pointers to arguments to the device (dp and hp might in
      fact have the same value on a unified-memory system).  */
-  if (mapnum > 0)
+  if (!params && mapnum > 0)
     {
       CUDA_CALL_ASSERT (cuMemcpyHtoDAsync, dp, (void *) hp,
 			mapnum * sizeof (void *), aq->cuda_stream);
@@ -1321,13 +1362,41 @@ GOMP_OFFLOAD_openacc_async_exec (void (*fn) (void *), size_t mapnum,
 	(struct nvptx_thread *) GOMP_PLUGIN_acc_thread ();
       block[1] = (void *) nvthd->ptx_dev;
     }
-  nvptx_exec (fn, mapnum, hostaddrs, devaddrs, dims, targ_mem_desc,
-	      dp, aq->cuda_stream);
 
-  if (mapnum > 0)
+  if (params)
+    nvptx_exec (fn, mapnum, hostaddrs, devaddrs, dims, targ_mem_desc,
+		hp, aq->cuda_stream);
+  else
+    {
+      void *kargs[1] = { &dp };
+      nvptx_exec (fn, mapnum, hostaddrs, devaddrs, dims, targ_mem_desc,
+		  kargs, aq->cuda_stream);
+    }
+
+  if (!params && mapnum > 0)
     GOMP_OFFLOAD_openacc_async_queue_callback (aq, cuda_free_argmem, block);
 }
 
+void
+GOMP_OFFLOAD_openacc_async_exec_params (void (*fn) (void *), size_t mapnum,
+				 void **hostaddrs, void **devaddrs,
+				 unsigned *dims, void *targ_mem_desc,
+				 struct goacc_asyncqueue *aq)
+{
+  openacc_async_exec_internal (fn, 1, mapnum, hostaddrs, devaddrs, dims,
+			       targ_mem_desc, aq);
+}
+
+void
+GOMP_OFFLOAD_openacc_async_exec (void (*fn) (void *), size_t mapnum,
+				 void **hostaddrs, void **devaddrs,
+				 unsigned *dims, void *targ_mem_desc,
+				 struct goacc_asyncqueue *aq)
+{
+  openacc_async_exec_internal (fn, 0, mapnum, hostaddrs, devaddrs, dims,
+			       targ_mem_desc, aq);
+}
+
 void *
 GOMP_OFFLOAD_openacc_create_thread_data (int ord)
 {
diff --git a/libgomp/target.c b/libgomp/target.c
index b34043f31423..91d0f25cde57 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -3580,6 +3580,7 @@ gomp_load_plugin_for_device (struct gomp_device_descr *device,
   if (device->capabilities & GOMP_OFFLOAD_CAP_OPENACC_200)
     {
       if (!DLSYM_OPT (openacc.exec, openacc_exec)
+	  || !DLSYM_OPT (openacc.exec_params, openacc_exec_params)
 	  || !DLSYM_OPT (openacc.create_thread_data,
 			 openacc_create_thread_data)
 	  || !DLSYM_OPT (openacc.destroy_thread_data,
@@ -3592,6 +3593,7 @@ gomp_load_plugin_for_device (struct gomp_device_descr *device,
 	  || !DLSYM_OPT (openacc.async.queue_callback,
 			 openacc_async_queue_callback)
 	  || !DLSYM_OPT (openacc.async.exec, openacc_async_exec)
+	  || !DLSYM_OPT (openacc.async.exec_params, openacc_async_exec_params)
 	  || !DLSYM_OPT (openacc.async.dev2host, openacc_async_dev2host)
 	  || !DLSYM_OPT (openacc.async.host2dev, openacc_async_host2dev))
 	{
diff --git a/libgomp/testsuite/Makefile.in b/libgomp/testsuite/Makefile.in
index 80315b15a7d5..db794f35d6de 100644
--- a/libgomp/testsuite/Makefile.in
+++ b/libgomp/testsuite/Makefile.in
@@ -168,6 +168,8 @@ INSTALL_SCRIPT = @INSTALL_SCRIPT@
 INSTALL_STRIP_PROGRAM = @INSTALL_STRIP_PROGRAM@
 LD = @LD@
 LDFLAGS = @LDFLAGS@
+LIBFFI = @LIBFFI@
+LIBFFIINCS = @LIBFFIINCS@
 LIBOBJS = @LIBOBJS@
 LIBS = @LIBS@
 LIBTOOL = @LIBTOOL@
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/combined-directives-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/combined-directives-1.c
index dad6d13eb608..c6abc1d724ab 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/combined-directives-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/combined-directives-1.c
@@ -1,6 +1,11 @@
 /* This test exercises combined directives.  */
 
+/* This test falls back to host execution because struct alias
+   analysis is deactivated on OpenACC parallel regions.  Consequently,
+   parloops can no longer disambiguate arrays a and b.  */
+
 /* { dg-do run } */
+/* { dg-xfail-if "n/a" { openacc_nvidia_accel_selected } { "-O2" } { "" } } */
 
 #include <stdlib.h>