Merge branch 'master' into devel/coarray_native.

Merge into devel/coarray_native to prepare for later merging of
coarray_native with master.

36 files changed, 1270 insertions, 451 deletions

diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index 0e37557..35b6689 100644
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,3 +1,172 @@
+2020-10-22  Jakub Jelinek  <jakub@redhat.com>
+
+	* testsuite/libgomp.c/target-41.c: New test.
+
+2020-10-22  Jakub Jelinek  <jakub@redhat.com>
+
+	* icv.c (omp_get_initial_device): Remove including corresponding
+	ialias.
+	* icv-device.c (omp_get_initial_device): New function.  Return
+	gomp_get_num_devices ().  Add ialias.
+	* target.c (resolve_device): Don't fail with
+	OMP_TARGET_OFFLOAD=mandatory if device_id is equal to
+	gomp_get_num_devices ().
+	(omp_target_alloc, omp_target_free, omp_target_is_present,
+	omp_target_memcpy, omp_target_memcpy_rect, omp_target_associate_ptr,
+	omp_target_disassociate_ptr, omp_pause_resource): Use
+	gomp_get_num_devices () instead of GOMP_DEVICE_HOST_FALLBACK on the
+	first use in the functions, in uses dominated by the
+	gomp_get_num_devices call use num_devices_openmp instead.
+	* libgomp.texi (omp_get_initial_device): Document.
+	* config/gcn/icv-device.c (omp_get_initial_device): New function.
+	Add ialias.
+	* config/nvptx/icv-device.c (omp_get_initial_device): Likewise.
+	* testsuite/libgomp.c/target-40.c: New test.
+
+2020-10-21  Jakub Jelinek  <jakub@redhat.com>
+
+	* env.c (parse_target_offload): Change new_offload var type to int,
+	preinitialize to -1, remove found var and test new_offload != -1
+	instead of found.
+
+2020-10-20  Jakub Jelinek  <jakub@redhat.com>
+
+	* target.c (gomp_target_init): Inside of the function, use automatic
+	variables corresponding to num_devices, num_devices_openmp and devices
+	global variables and update the globals only at the end of the
+	function.
+
+2020-10-20  Kwok Cheung Yeung  <kcy@codesourcery.com>
+
+	* env.c (gomp_target_offload_var): New.
+	(parse_target_offload): New.
+	(handle_omp_display_env): Print value of OMP_TARGET_OFFLOAD.
+	(initialize_env): Parse OMP_TARGET_OFFLOAD.
+	* libgomp.h (gomp_target_offload_t): New.
+	(gomp_target_offload_var): New.
+	* libgomp.texi (OMP_TARGET_OFFLOAD): New section.
+	* target.c (resolve_device): Generate error if device not found and
+	offloading is mandatory.
+	(gomp_target_fallback): Generate error if offloading is mandatory.
+	(GOMP_target): Add argument in call to gomp_target_fallback.
+	(GOMP_target_ext): Likewise.
+	(gomp_target_data_fallback): Generate error if offloading is mandatory.
+	(GOMP_target_data): Add argument in call to gomp_target_data_fallback.
+	(GOMP_target_data_ext): Likewise.
+	(gomp_target_task_fn): Add argument in call to gomp_target_fallback.
+	(gomp_target_init): Return early if offloading is disabled.
+
+2020-10-15  Kwok Cheung Yeung  <kcy@codesourcery.com>
+
+	* libgomp.texi (omp_get_max_active_levels): Modify description.
+	(omp_get_supported_active_levels): Make descriptions consistent.
+
+2020-10-14  Jakub Jelinek  <jakub@redhat.com>
+
+	* libgomp.texi (omp_get_supported_active_levels): Fix a typo.
+
+2020-10-13  Kwok Cheung Yeung  <kcy@codesourcery.com>
+
+	* env.c (gomp_max_active_levels_var): Initialize to
+	gomp_supported_active_levels.
+	(initialize_env): Limit gomp_max_active_levels_var to be at most
+	equal to gomp_supported_active_levels.
+	* fortran.c (omp_get_supported_active_levels): Add ialias_redirect.
+	(omp_get_supported_active_levels_): New.
+	* icv.c (omp_set_max_active_levels): Limit gomp_max_active_levels_var
+	to at most equal to gomp_supported_active_levels.
+	(omp_get_supported_active_levels): New.
+	* libgomp.h (gomp_supported_active_levels): New.
+	* libgomp.map (OMP_5.0.1): Add omp_get_supported_active_levels and
+	omp_get_supported_active_levels_.
+	* libgomp.texi (omp_get_supported_active_levels): New.
+	(omp_set_max_active_levels): Update.  Add reference to
+	omp_get_supported_active_levels.
+	* omp.h.in (omp_get_supported_active_levels): New.
+	* omp_lib.f90.in (omp_get_supported_active_levels): New.
+	* omp_lib.h.in (omp_get_supported_active_levels): New.
+	* testsuite/libgomp.c/lib-2.c (main): Check omp_get_max_active_levels
+	against omp_get_supported_active_levels.
+	* testsuite/libgomp.fortran/lib4.f90 (lib4): Likewise.
+
+2020-10-11  Clement Chigot  <clement.chigot@atos.net>
+
+	* config/t-aix: Delete and recreate libgomp before creating
+	FAT library.
+
+2020-10-08  Tom de Vries  <tdevries@suse.de>
+
+	PR libgomp/81802
+	* plugin/plugin-nvptx.c (GOMP_OFFLOAD_run): Report launch
+	dimensions.
+
+2020-10-06  Tom de Vries  <tdevries@suse.de>
+
+	* testsuite/libgomp.oacc-fortran/declare-5.f90: Add xfail for PR92790.
+
+2020-10-06  Tom de Vries  <tdevries@suse.de>
+
+	PR middle-end/90861
+	* testsuite/libgomp.oacc-c-c++-common/declare-vla.c: Remove xfail.
+
+2020-10-05  Tom de Vries  <tdevries@suse.de>
+
+	PR fortran/95654
+	* testsuite/libgomp.fortran/pr95654.f90: New test.
+
+2020-10-02  Tobias Burnus  <tobias@codesourcery.com>
+
+	* Makefile.in: Regenerate with automake 1.15.1.
+	* aclocal.m4: Likewise.
+	* configure: Likewise.
+	* testsuite/Makefile.in: Likewise.
+
+2020-09-30  Andrew Stubbs  <ams@codesourcery.com>
+
+	* parallel.c (gomp_resolve_num_threads): Ignore nest_var on nvptx
+	and amdgcn targets.
+
+2020-09-30  Tobias Burnus  <tobias@codesourcery.com>
+
+	* testsuite/libgomp.fortran/declare-target-3.f90: New test.
+
+2020-09-29  Andrew Stubbs  <ams@codesourcery.com>
+
+	* config/gcn/bar.c (gomp_barrier_wait_end): Skip the barrier if the
+	total number of threads is one.
+	(gomp_team_barrier_wake): Likewise.
+	(gomp_team_barrier_wait_end): Likewise.
+	(gomp_team_barrier_wait_cancel_end): Likewise.
+	* config/nvptx/bar.c (gomp_barrier_wait_end): Likewise.
+	(gomp_team_barrier_wake): Likewise.
+	(gomp_team_barrier_wait_end): Likewise.
+	(gomp_team_barrier_wait_cancel_end): Likewise.
+	* testsuite/libgomp.c-c++-common/nested-parallel-unbalanced.c: New test.
+
+2020-09-28  Tobias Burnus  <tobias@codesourcery.com>
+
+	PR middle-end/96390
+	* testsuite/libgomp.c++/pr96390.C: New test.
+	* testsuite/libgomp.c-c++-common/pr96390.c: New test.
+
+2020-09-27  Clement Chigot  <clement.chigot@atos.net>
+
+	* config/t-aix: Use $(AR) without -X32_64.
+
+2020-09-25  Jakub Jelinek  <jakub@redhat.com>
+
+	* testsuite/libgomp.c/loop-25.c: New test.
+
+2020-09-22  Tobias Burnus  <tobias@codesourcery.com>
+
+	PR fortran/95654
+	* testsuite/libgomp.fortran/pr66199-5.f90: Make stop codes unique.
+
+2020-09-22  Tom de Vries  <tdevries@suse.de>
+
+	* plugin/plugin-nvptx.c (link_ptx): Print elog if cuLinkComplete call
+	fails.
+
 2020-09-16  Nathan Sidwell  <nathan@acm.org>
 
 	* testsuite/libgomp.c++/udr-3.C: Add missing ctor.
diff --git a/libgomp/Makefile.in b/libgomp/Makefile.in
index 2dc2168..00d5e29 100644
--- a/libgomp/Makefile.in
+++ b/libgomp/Makefile.in
@@ -1,7 +1,7 @@
-# Makefile.in generated by automake 1.16.1 from Makefile.am.
+# Makefile.in generated by automake 1.15.1 from Makefile.am.
 # @configure_input@
 
-# Copyright (C) 1994-2018 Free Software Foundation, Inc.
+# Copyright (C) 1994-2017 Free Software Foundation, Inc.
 
 # This Makefile.in is free software; the Free Software Foundation
 # gives unlimited permission to copy and/or distribute it,
@@ -236,32 +236,7 @@ am__v_at_0 = @
 am__v_at_1 = 
 DEFAULT_INCLUDES = -I.@am__isrc@
 depcomp = $(SHELL) $(top_srcdir)/../depcomp
-am__maybe_remake_depfiles = depfiles
-am__depfiles_remade = ./$(DEPDIR)/affinity-fmt.Plo \
-	./$(DEPDIR)/affinity.Plo ./$(DEPDIR)/alloc.Plo \
-	./$(DEPDIR)/allocator.Plo ./$(DEPDIR)/atomic.Plo \
-	./$(DEPDIR)/bar.Plo ./$(DEPDIR)/barrier.Plo \
-	./$(DEPDIR)/critical.Plo ./$(DEPDIR)/env.Plo \
-	./$(DEPDIR)/error.Plo ./$(DEPDIR)/fortran.Plo \
-	./$(DEPDIR)/icv-device.Plo ./$(DEPDIR)/icv.Plo \
-	./$(DEPDIR)/iter.Plo ./$(DEPDIR)/iter_ull.Plo \
-	./$(DEPDIR)/libgomp-plugin.Plo \
-	./$(DEPDIR)/libgomp_plugin_gcn_la-plugin-gcn.Plo \
-	./$(DEPDIR)/libgomp_plugin_nvptx_la-plugin-nvptx.Plo \
-	./$(DEPDIR)/lock.Plo ./$(DEPDIR)/loop.Plo \
-	./$(DEPDIR)/loop_ull.Plo ./$(DEPDIR)/mutex.Plo \
-	./$(DEPDIR)/oacc-async.Plo ./$(DEPDIR)/oacc-cuda.Plo \
-	./$(DEPDIR)/oacc-host.Plo ./$(DEPDIR)/oacc-init.Plo \
-	./$(DEPDIR)/oacc-mem.Plo ./$(DEPDIR)/oacc-parallel.Plo \
-	./$(DEPDIR)/oacc-plugin.Plo ./$(DEPDIR)/oacc-profiling.Plo \
-	./$(DEPDIR)/oacc-target.Plo ./$(DEPDIR)/ordered.Plo \
-	./$(DEPDIR)/parallel.Plo ./$(DEPDIR)/priority_queue.Plo \
-	./$(DEPDIR)/proc.Plo ./$(DEPDIR)/ptrlock.Plo \
-	./$(DEPDIR)/sections.Plo ./$(DEPDIR)/sem.Plo \
-	./$(DEPDIR)/single.Plo ./$(DEPDIR)/splay-tree.Plo \
-	./$(DEPDIR)/target.Plo ./$(DEPDIR)/task.Plo \
-	./$(DEPDIR)/team.Plo ./$(DEPDIR)/teams.Plo \
-	./$(DEPDIR)/time.Plo ./$(DEPDIR)/work.Plo
+am__depfiles_maybe = depfiles
 am__mv = mv -f
 COMPILE = $(CC) $(DEFS) $(DEFAULT_INCLUDES) $(INCLUDES) $(AM_CPPFLAGS) \
 	$(CPPFLAGS) $(AM_CFLAGS) $(CFLAGS)
@@ -669,8 +644,8 @@ Makefile: $(srcdir)/Makefile.in $(top_builddir)/config.status
 	    echo ' $(SHELL) ./config.status'; \
 	    $(SHELL) ./config.status;; \
 	  *) \
-	    echo ' cd $(top_builddir) && $(SHELL) ./config.status $@ $(am__maybe_remake_depfiles)'; \
-	    cd $(top_builddir) && $(SHELL) ./config.status $@ $(am__maybe_remake_depfiles);; \
+	    echo ' cd $(top_builddir) && $(SHELL) ./config.status $@ $(am__depfiles_maybe)'; \
+	    cd $(top_builddir) && $(SHELL) ./config.status $@ $(am__depfiles_maybe);; \
 	esac;
 $(top_srcdir)/plugin/Makefrag.am $(top_srcdir)/../multilib.am $(am__empty):
 
@@ -758,58 +733,52 @@ mostlyclean-compile:
 distclean-compile:
 	-rm -f *.tab.c
 
-@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/affinity-fmt.Plo@am__quote@ # am--include-marker
-@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/affinity.Plo@am__quote@ # am--include-marker
-@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/alloc.Plo@am__quote@ # am--include-marker
-@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/allocator.Plo@am__quote@ # am--include-marker
-@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/atomic.Plo@am__quote@ # am--include-marker
-@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/bar.Plo@am__quote@ # am--include-marker
-@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/barrier.Plo@am__quote@ # am--include-marker
-@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/critical.Plo@am__quote@ # am--include-marker
-@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/env.Plo@am__quote@ # am--include-marker
-@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/error.Plo@am__quote@ # am--include-marker
-@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/fortran.Plo@am__quote@ # am--include-marker
-@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/icv-device.Plo@am__quote@ # am--include-marker
-@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/icv.Plo@am__quote@ # am--include-marker
-@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/iter.Plo@am__quote@ # am--include-marker
-@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/iter_ull.Plo@am__quote@ # am--include-marker
-@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/libgomp-plugin.Plo@am__quote@ # am--include-marker
-@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/libgomp_plugin_gcn_la-plugin-gcn.Plo@am__quote@ # am--include-marker
-@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/libgomp_plugin_nvptx_la-plugin-nvptx.Plo@am__quote@ # am--include-marker
-@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/lock.Plo@am__quote@ # am--include-marker
-@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/loop.Plo@am__quote@ # am--include-marker
-@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/loop_ull.Plo@am__quote@ # am--include-marker
-@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/mutex.Plo@am__quote@ # am--include-marker
-@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/oacc-async.Plo@am__quote@ # am--include-marker
-@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/oacc-cuda.Plo@am__quote@ # am--include-marker
-@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/oacc-host.Plo@am__quote@ # am--include-marker
-@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/oacc-init.Plo@am__quote@ # am--include-marker
-@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/oacc-mem.Plo@am__quote@ # am--include-marker
-@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/oacc-parallel.Plo@am__quote@ # am--include-marker
-@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/oacc-plugin.Plo@am__quote@ # am--include-marker
-@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/oacc-profiling.Plo@am__quote@ # am--include-marker
-@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/oacc-target.Plo@am__quote@ # am--include-marker
-@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/ordered.Plo@am__quote@ # am--include-marker
-@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/parallel.Plo@am__quote@ # am--include-marker
-@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/priority_queue.Plo@am__quote@ # am--include-marker
-@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/proc.Plo@am__quote@ # am--include-marker
-@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/ptrlock.Plo@am__quote@ # am--include-marker
-@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/sections.Plo@am__quote@ # am--include-marker
-@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/sem.Plo@am__quote@ # am--include-marker
-@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/single.Plo@am__quote@ # am--include-marker
-@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/splay-tree.Plo@am__quote@ # am--include-marker
-@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/target.Plo@am__quote@ # am--include-marker
-@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/task.Plo@am__quote@ # am--include-marker
-@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/team.Plo@am__quote@ # am--include-marker
-@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/teams.Plo@am__quote@ # am--include-marker
-@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/time.Plo@am__quote@ # am--include-marker
-@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/work.Plo@am__quote@ # am--include-marker
-
-$(am__depfiles_remade):
-	@$(MKDIR_P) $(@D)
-	@echo '# dummy' >$@-t && $(am__mv) $@-t $@
-
-am--depfiles: $(am__depfiles_remade)
+@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/affinity-fmt.Plo@am__quote@
+@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/affinity.Plo@am__quote@
+@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/alloc.Plo@am__quote@
+@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/allocator.Plo@am__quote@
+@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/atomic.Plo@am__quote@
+@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/bar.Plo@am__quote@
+@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/barrier.Plo@am__quote@
+@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/critical.Plo@am__quote@
+@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/env.Plo@am__quote@
+@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/error.Plo@am__quote@
+@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/fortran.Plo@am__quote@
+@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/icv-device.Plo@am__quote@
+@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/icv.Plo@am__quote@
+@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/iter.Plo@am__quote@
+@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/iter_ull.Plo@am__quote@
+@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/libgomp-plugin.Plo@am__quote@
+@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/libgomp_plugin_gcn_la-plugin-gcn.Plo@am__quote@
+@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/libgomp_plugin_nvptx_la-plugin-nvptx.Plo@am__quote@
+@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/lock.Plo@am__quote@
+@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/loop.Plo@am__quote@
+@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/loop_ull.Plo@am__quote@
+@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/mutex.Plo@am__quote@
+@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/oacc-async.Plo@am__quote@
+@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/oacc-cuda.Plo@am__quote@
+@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/oacc-host.Plo@am__quote@
+@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/oacc-init.Plo@am__quote@
+@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/oacc-mem.Plo@am__quote@
+@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/oacc-parallel.Plo@am__quote@
+@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/oacc-plugin.Plo@am__quote@
+@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/oacc-profiling.Plo@am__quote@
+@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/oacc-target.Plo@am__quote@
+@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/ordered.Plo@am__quote@
+@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/parallel.Plo@am__quote@
+@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/priority_queue.Plo@am__quote@
+@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/proc.Plo@am__quote@
+@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/ptrlock.Plo@am__quote@
+@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/sections.Plo@am__quote@
+@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/sem.Plo@am__quote@
+@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/single.Plo@am__quote@
+@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/splay-tree.Plo@am__quote@
+@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/target.Plo@am__quote@
+@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/task.Plo@am__quote@
+@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/team.Plo@am__quote@
+@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/teams.Plo@am__quote@
+@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/time.Plo@am__quote@
+@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/work.Plo@am__quote@
 
 .c.o:
 @am__fastdepCC_TRUE@	$(AM_V_CC)$(COMPILE) -MT $@ -MD -MP -MF $(DEPDIR)/$*.Tpo -c -o $@ $<
@@ -1195,52 +1164,7 @@ clean-am: clean-aminfo clean-generic clean-libtool clean-local \
 
 distclean: distclean-recursive
 	-rm -f $(am__CONFIG_DISTCLEAN_FILES)
-		-rm -f ./$(DEPDIR)/affinity-fmt.Plo
-	-rm -f ./$(DEPDIR)/affinity.Plo
-	-rm -f ./$(DEPDIR)/alloc.Plo
-	-rm -f ./$(DEPDIR)/allocator.Plo
-	-rm -f ./$(DEPDIR)/atomic.Plo
-	-rm -f ./$(DEPDIR)/bar.Plo
-	-rm -f ./$(DEPDIR)/barrier.Plo
-	-rm -f ./$(DEPDIR)/critical.Plo
-	-rm -f ./$(DEPDIR)/env.Plo
-	-rm -f ./$(DEPDIR)/error.Plo
-	-rm -f ./$(DEPDIR)/fortran.Plo
-	-rm -f ./$(DEPDIR)/icv-device.Plo
-	-rm -f ./$(DEPDIR)/icv.Plo
-	-rm -f ./$(DEPDIR)/iter.Plo
-	-rm -f ./$(DEPDIR)/iter_ull.Plo
-	-rm -f ./$(DEPDIR)/libgomp-plugin.Plo
-	-rm -f ./$(DEPDIR)/libgomp_plugin_gcn_la-plugin-gcn.Plo
-	-rm -f ./$(DEPDIR)/libgomp_plugin_nvptx_la-plugin-nvptx.Plo
-	-rm -f ./$(DEPDIR)/lock.Plo
-	-rm -f ./$(DEPDIR)/loop.Plo
-	-rm -f ./$(DEPDIR)/loop_ull.Plo
-	-rm -f ./$(DEPDIR)/mutex.Plo
-	-rm -f ./$(DEPDIR)/oacc-async.Plo
-	-rm -f ./$(DEPDIR)/oacc-cuda.Plo
-	-rm -f ./$(DEPDIR)/oacc-host.Plo
-	-rm -f ./$(DEPDIR)/oacc-init.Plo
-	-rm -f ./$(DEPDIR)/oacc-mem.Plo
-	-rm -f ./$(DEPDIR)/oacc-parallel.Plo
-	-rm -f ./$(DEPDIR)/oacc-plugin.Plo
-	-rm -f ./$(DEPDIR)/oacc-profiling.Plo
-	-rm -f ./$(DEPDIR)/oacc-target.Plo
-	-rm -f ./$(DEPDIR)/ordered.Plo
-	-rm -f ./$(DEPDIR)/parallel.Plo
-	-rm -f ./$(DEPDIR)/priority_queue.Plo
-	-rm -f ./$(DEPDIR)/proc.Plo
-	-rm -f ./$(DEPDIR)/ptrlock.Plo
-	-rm -f ./$(DEPDIR)/sections.Plo
-	-rm -f ./$(DEPDIR)/sem.Plo
-	-rm -f ./$(DEPDIR)/single.Plo
-	-rm -f ./$(DEPDIR)/splay-tree.Plo
-	-rm -f ./$(DEPDIR)/target.Plo
-	-rm -f ./$(DEPDIR)/task.Plo
-	-rm -f ./$(DEPDIR)/team.Plo
-	-rm -f ./$(DEPDIR)/teams.Plo
-	-rm -f ./$(DEPDIR)/time.Plo
-	-rm -f ./$(DEPDIR)/work.Plo
+	-rm -rf ./$(DEPDIR)
 	-rm -f Makefile
 distclean-am: clean-am distclean-compile distclean-generic \
 	distclean-hdr distclean-libtool distclean-local distclean-tags
@@ -1381,52 +1305,7 @@ installcheck-am:
 maintainer-clean: maintainer-clean-recursive
 	-rm -f $(am__CONFIG_DISTCLEAN_FILES)
 	-rm -rf $(top_srcdir)/autom4te.cache
-		-rm -f ./$(DEPDIR)/affinity-fmt.Plo
-	-rm -f ./$(DEPDIR)/affinity.Plo
-	-rm -f ./$(DEPDIR)/alloc.Plo
-	-rm -f ./$(DEPDIR)/allocator.Plo
-	-rm -f ./$(DEPDIR)/atomic.Plo
-	-rm -f ./$(DEPDIR)/bar.Plo
-	-rm -f ./$(DEPDIR)/barrier.Plo
-	-rm -f ./$(DEPDIR)/critical.Plo
-	-rm -f ./$(DEPDIR)/env.Plo
-	-rm -f ./$(DEPDIR)/error.Plo
-	-rm -f ./$(DEPDIR)/fortran.Plo
-	-rm -f ./$(DEPDIR)/icv-device.Plo
-	-rm -f ./$(DEPDIR)/icv.Plo
-	-rm -f ./$(DEPDIR)/iter.Plo
-	-rm -f ./$(DEPDIR)/iter_ull.Plo
-	-rm -f ./$(DEPDIR)/libgomp-plugin.Plo
-	-rm -f ./$(DEPDIR)/libgomp_plugin_gcn_la-plugin-gcn.Plo
-	-rm -f ./$(DEPDIR)/libgomp_plugin_nvptx_la-plugin-nvptx.Plo
-	-rm -f ./$(DEPDIR)/lock.Plo
-	-rm -f ./$(DEPDIR)/loop.Plo
-	-rm -f ./$(DEPDIR)/loop_ull.Plo
-	-rm -f ./$(DEPDIR)/mutex.Plo
-	-rm -f ./$(DEPDIR)/oacc-async.Plo
-	-rm -f ./$(DEPDIR)/oacc-cuda.Plo
-	-rm -f ./$(DEPDIR)/oacc-host.Plo
-	-rm -f ./$(DEPDIR)/oacc-init.Plo
-	-rm -f ./$(DEPDIR)/oacc-mem.Plo
-	-rm -f ./$(DEPDIR)/oacc-parallel.Plo
-	-rm -f ./$(DEPDIR)/oacc-plugin.Plo
-	-rm -f ./$(DEPDIR)/oacc-profiling.Plo
-	-rm -f ./$(DEPDIR)/oacc-target.Plo
-	-rm -f ./$(DEPDIR)/ordered.Plo
-	-rm -f ./$(DEPDIR)/parallel.Plo
-	-rm -f ./$(DEPDIR)/priority_queue.Plo
-	-rm -f ./$(DEPDIR)/proc.Plo
-	-rm -f ./$(DEPDIR)/ptrlock.Plo
-	-rm -f ./$(DEPDIR)/sections.Plo
-	-rm -f ./$(DEPDIR)/sem.Plo
-	-rm -f ./$(DEPDIR)/single.Plo
-	-rm -f ./$(DEPDIR)/splay-tree.Plo
-	-rm -f ./$(DEPDIR)/target.Plo
-	-rm -f ./$(DEPDIR)/task.Plo
-	-rm -f ./$(DEPDIR)/team.Plo
-	-rm -f ./$(DEPDIR)/teams.Plo
-	-rm -f ./$(DEPDIR)/time.Plo
-	-rm -f ./$(DEPDIR)/work.Plo
+	-rm -rf ./$(DEPDIR)
 	-rm -f Makefile
 maintainer-clean-am: distclean-am maintainer-clean-aminfo \
 	maintainer-clean-generic maintainer-clean-local
@@ -1453,8 +1332,8 @@ uninstall-am: uninstall-dvi-am uninstall-html-am uninstall-info-am \
 .MAKE: $(am__recursive_targets) all install-am install-strip
 
 .PHONY: $(am__recursive_targets) CTAGS GTAGS TAGS all all-am all-local \
-	am--depfiles am--refresh check check-am clean clean-aminfo \
-	clean-cscope clean-generic clean-libtool clean-local \
+	am--refresh check check-am clean clean-aminfo clean-cscope \
+	clean-generic clean-libtool clean-local \
 	clean-toolexeclibLTLIBRARIES cscope cscopelist-am ctags \
 	ctags-am dist-info distclean distclean-compile \
 	distclean-generic distclean-hdr distclean-libtool \
diff --git a/libgomp/aclocal.m4 b/libgomp/aclocal.m4
index 471963b..55d9d71 100644
--- a/libgomp/aclocal.m4
+++ b/libgomp/aclocal.m4
@@ -1,6 +1,6 @@
-# generated automatically by aclocal 1.16.1 -*- Autoconf -*-
+# generated automatically by aclocal 1.15.1 -*- Autoconf -*-
 
-# Copyright (C) 1996-2018 Free Software Foundation, Inc.
+# Copyright (C) 1996-2017 Free Software Foundation, Inc.
 
 # This file is free software; the Free Software Foundation
 # gives unlimited permission to copy and/or distribute it,
@@ -20,7 +20,7 @@ You have another version of autoconf.  It may work, but is not guaranteed to.
 If you have problems, you may need to regenerate the build system entirely.
 To do so, use the procedure documented by the package, typically 'autoreconf'.])])
 
-# Copyright (C) 2002-2018 Free Software Foundation, Inc.
+# Copyright (C) 2002-2017 Free Software Foundation, Inc.
 #
 # This file is free software; the Free Software Foundation
 # gives unlimited permission to copy and/or distribute it,
@@ -32,10 +32,10 @@ To do so, use the procedure documented by the package, typically 'autoreconf'.])
 # generated from the m4 files accompanying Automake X.Y.
 # (This private macro should not be called outside this file.)
 AC_DEFUN([AM_AUTOMAKE_VERSION],
-[am__api_version='1.16'
+[am__api_version='1.15'
 dnl Some users find AM_AUTOMAKE_VERSION and mistake it for a way to
 dnl require some minimum version.  Point them to the right macro.
-m4_if([$1], [1.16.1], [],
+m4_if([$1], [1.15.1], [],
       [AC_FATAL([Do not call $0, use AM_INIT_AUTOMAKE([$1]).])])dnl
 ])
 
@@ -51,14 +51,14 @@ m4_define([_AM_AUTOCONF_VERSION], [])
 # Call AM_AUTOMAKE_VERSION and AM_AUTOMAKE_VERSION so they can be traced.
 # This function is AC_REQUIREd by AM_INIT_AUTOMAKE.
 AC_DEFUN([AM_SET_CURRENT_AUTOMAKE_VERSION],
-[AM_AUTOMAKE_VERSION([1.16.1])dnl
+[AM_AUTOMAKE_VERSION([1.15.1])dnl
 m4_ifndef([AC_AUTOCONF_VERSION],
   [m4_copy([m4_PACKAGE_VERSION], [AC_AUTOCONF_VERSION])])dnl
 _AM_AUTOCONF_VERSION(m4_defn([AC_AUTOCONF_VERSION]))])
 
 # AM_AUX_DIR_EXPAND                                         -*- Autoconf -*-
 
-# Copyright (C) 2001-2018 Free Software Foundation, Inc.
+# Copyright (C) 2001-2017 Free Software Foundation, Inc.
 #
 # This file is free software; the Free Software Foundation
 # gives unlimited permission to copy and/or distribute it,
@@ -110,7 +110,7 @@ am_aux_dir=`cd "$ac_aux_dir" && pwd`
 
 # AM_CONDITIONAL                                            -*- Autoconf -*-
 
-# Copyright (C) 1997-2018 Free Software Foundation, Inc.
+# Copyright (C) 1997-2017 Free Software Foundation, Inc.
 #
 # This file is free software; the Free Software Foundation
 # gives unlimited permission to copy and/or distribute it,
@@ -141,7 +141,7 @@ AC_CONFIG_COMMANDS_PRE(
 Usually this means the macro was only invoked conditionally.]])
 fi])])
 
-# Copyright (C) 1999-2018 Free Software Foundation, Inc.
+# Copyright (C) 1999-2017 Free Software Foundation, Inc.
 #
 # This file is free software; the Free Software Foundation
 # gives unlimited permission to copy and/or distribute it,
@@ -332,12 +332,13 @@ _AM_SUBST_NOTMAKE([am__nodep])dnl
 
 # Generate code to set up dependency tracking.              -*- Autoconf -*-
 
-# Copyright (C) 1999-2018 Free Software Foundation, Inc.
+# Copyright (C) 1999-2017 Free Software Foundation, Inc.
 #
 # This file is free software; the Free Software Foundation
 # gives unlimited permission to copy and/or distribute it,
 # with or without modifications, as long as this notice is preserved.
 
+
 # _AM_OUTPUT_DEPENDENCY_COMMANDS
 # ------------------------------
 AC_DEFUN([_AM_OUTPUT_DEPENDENCY_COMMANDS],
@@ -345,41 +346,49 @@ AC_DEFUN([_AM_OUTPUT_DEPENDENCY_COMMANDS],
   # Older Autoconf quotes --file arguments for eval, but not when files
   # are listed without --file.  Let's play safe and only enable the eval
   # if we detect the quoting.
-  # TODO: see whether this extra hack can be removed once we start
-  # requiring Autoconf 2.70 or later.
-  AS_CASE([$CONFIG_FILES],
-          [*\'*], [eval set x "$CONFIG_FILES"],
-          [*], [set x $CONFIG_FILES])
+  case $CONFIG_FILES in
+  *\'*) eval set x "$CONFIG_FILES" ;;
+  *)   set x $CONFIG_FILES ;;
+  esac
   shift
-  # Used to flag and report bootstrapping failures.
-  am_rc=0
-  for am_mf
+  for mf
   do
     # Strip MF so we end up with the name of the file.
-    am_mf=`AS_ECHO(["$am_mf"]) | sed -e 's/:.*$//'`
-    # Check whether this is an Automake generated Makefile which includes
-    # dependency-tracking related rules and includes.
-    # Grep'ing the whole file directly is not great: AIX grep has a line
+    mf=`echo "$mf" | sed -e 's/:.*$//'`
+    # Check whether this is an Automake generated Makefile or not.
+    # We used to match only the files named 'Makefile.in', but
+    # some people rename them; so instead we look at the file content.
+    # Grep'ing the first line is not enough: some people post-process
+    # each Makefile.in and add a new line on top of each file to say so.
+    # Grep'ing the whole file is not good either: AIX grep has a line
     # limit of 2048, but all sed's we know have understand at least 4000.
-    sed -n 's,^am--depfiles:.*,X,p' "$am_mf" | grep X >/dev/null 2>&1 \
-      || continue
-    am_dirpart=`AS_DIRNAME(["$am_mf"])`
-    am_filepart=`AS_BASENAME(["$am_mf"])`
-    AM_RUN_LOG([cd "$am_dirpart" \
-      && sed -e '/# am--include-marker/d' "$am_filepart" \
-        | $MAKE -f - am--depfiles]) || am_rc=$?
+    if sed -n 's,^#.*generated by automake.*,X,p' "$mf" | grep X >/dev/null 2>&1; then
+      dirpart=`AS_DIRNAME("$mf")`
+    else
+      continue
+    fi
+    # Extract the definition of DEPDIR, am__include, and am__quote
+    # from the Makefile without running 'make'.
+    DEPDIR=`sed -n 's/^DEPDIR = //p' < "$mf"`
+    test -z "$DEPDIR" && continue
+    am__include=`sed -n 's/^am__include = //p' < "$mf"`
+    test -z "$am__include" && continue
+    am__quote=`sed -n 's/^am__quote = //p' < "$mf"`
+    # Find all dependency output files, they are included files with
+    # $(DEPDIR) in their names.  We invoke sed twice because it is the
+    # simplest approach to changing $(DEPDIR) to its actual value in the
+    # expansion.
+    for file in `sed -n "
+      s/^$am__include $am__quote\(.*(DEPDIR).*\)$am__quote"'$/\1/p' <"$mf" | \
+	 sed -e 's/\$(DEPDIR)/'"$DEPDIR"'/g'`; do
+      # Make sure the directory exists.
+      test -f "$dirpart/$file" && continue
+      fdir=`AS_DIRNAME(["$file"])`
+      AS_MKDIR_P([$dirpart/$fdir])
+      # echo "creating $dirpart/$file"
+      echo '# dummy' > "$dirpart/$file"
+    done
   done
-  if test $am_rc -ne 0; then
-    AC_MSG_FAILURE([Something went wrong bootstrapping makefile fragments
-    for automatic dependency tracking.  Try re-running configure with the
-    '--disable-dependency-tracking' option to at least be able to build
-    the package (albeit without support for automatic dependency tracking).])
-  fi
-  AS_UNSET([am_dirpart])
-  AS_UNSET([am_filepart])
-  AS_UNSET([am_mf])
-  AS_UNSET([am_rc])
-  rm -f conftest-deps.mk
 }
 ])# _AM_OUTPUT_DEPENDENCY_COMMANDS
 
@@ -388,17 +397,18 @@ AC_DEFUN([_AM_OUTPUT_DEPENDENCY_COMMANDS],
 # -----------------------------
 # This macro should only be invoked once -- use via AC_REQUIRE.
 #
-# This code is only required when automatic dependency tracking is enabled.
-# This creates each '.Po' and '.Plo' makefile fragment that we'll need in
-# order to bootstrap the dependency handling code.
+# This code is only required when automatic dependency tracking
+# is enabled.  FIXME.  This creates each '.P' file that we will
+# need in order to bootstrap the dependency handling code.
 AC_DEFUN([AM_OUTPUT_DEPENDENCY_COMMANDS],
 [AC_CONFIG_COMMANDS([depfiles],
      [test x"$AMDEP_TRUE" != x"" || _AM_OUTPUT_DEPENDENCY_COMMANDS],
-     [AMDEP_TRUE="$AMDEP_TRUE" MAKE="${MAKE-make}"])])
+     [AMDEP_TRUE="$AMDEP_TRUE" ac_aux_dir="$ac_aux_dir"])
+])
 
 # Do all the work for Automake.                             -*- Autoconf -*-
 
-# Copyright (C) 1996-2018 Free Software Foundation, Inc.
+# Copyright (C) 1996-2017 Free Software Foundation, Inc.
 #
 # This file is free software; the Free Software Foundation
 # gives unlimited permission to copy and/or distribute it,
@@ -485,8 +495,8 @@ AC_REQUIRE([AM_PROG_INSTALL_STRIP])dnl
 AC_REQUIRE([AC_PROG_MKDIR_P])dnl
 # For better backward compatibility.  To be removed once Automake 1.9.x
 # dies out for good.  For more background, see:
-# <https://lists.gnu.org/archive/html/automake/2012-07/msg00001.html>
-# <https://lists.gnu.org/archive/html/automake/2012-07/msg00014.html>
+# <http://lists.gnu.org/archive/html/automake/2012-07/msg00001.html>
+# <http://lists.gnu.org/archive/html/automake/2012-07/msg00014.html>
 AC_SUBST([mkdir_p], ['$(MKDIR_P)'])
 # We need awk for the "check" target (and possibly the TAP driver).  The
 # system "awk" is bad on some platforms.
@@ -553,7 +563,7 @@ END
 Aborting the configuration process, to ensure you take notice of the issue.
 
 You can download and install GNU coreutils to get an 'rm' implementation
-that behaves properly: <https://www.gnu.org/software/coreutils/>.
+that behaves properly: <http://www.gnu.org/software/coreutils/>.
 
 If you want to complete the configuration process using your problematic
 'rm' anyway, export the environment variable ACCEPT_INFERIOR_RM_PROGRAM
@@ -595,7 +605,7 @@ for _am_header in $config_headers :; do
 done
 echo "timestamp for $_am_arg" >`AS_DIRNAME(["$_am_arg"])`/stamp-h[]$_am_stamp_count])
 
-# Copyright (C) 2001-2018 Free Software Foundation, Inc.
+# Copyright (C) 2001-2017 Free Software Foundation, Inc.
 #
 # This file is free software; the Free Software Foundation
 # gives unlimited permission to copy and/or distribute it,
@@ -619,7 +629,7 @@ AC_SUBST([install_sh])])
 # Add --enable-maintainer-mode option to configure.         -*- Autoconf -*-
 # From Jim Meyering
 
-# Copyright (C) 1996-2018 Free Software Foundation, Inc.
+# Copyright (C) 1996-2017 Free Software Foundation, Inc.
 #
 # This file is free software; the Free Software Foundation
 # gives unlimited permission to copy and/or distribute it,
@@ -654,7 +664,7 @@ AC_MSG_CHECKING([whether to enable maintainer-specific portions of Makefiles])
 
 # Check to see how 'make' treats includes.	            -*- Autoconf -*-
 
-# Copyright (C) 2001-2018 Free Software Foundation, Inc.
+# Copyright (C) 2001-2017 Free Software Foundation, Inc.
 #
 # This file is free software; the Free Software Foundation
 # gives unlimited permission to copy and/or distribute it,
@@ -662,42 +672,49 @@ AC_MSG_CHECKING([whether to enable maintainer-specific portions of Makefiles])
 
 # AM_MAKE_INCLUDE()
 # -----------------
-# Check whether make has an 'include' directive that can support all
-# the idioms we need for our automatic dependency tracking code.
+# Check to see how make treats includes.
 AC_DEFUN([AM_MAKE_INCLUDE],
-[AC_MSG_CHECKING([whether ${MAKE-make} supports the include directive])
-cat > confinc.mk << 'END'
+[am_make=${MAKE-make}
+cat > confinc << 'END'
 am__doit:
-	@echo this is the am__doit target >confinc.out
+	@echo this is the am__doit target
 .PHONY: am__doit
 END
+# If we don't find an include directive, just comment out the code.
+AC_MSG_CHECKING([for style of include used by $am_make])
 am__include="#"
 am__quote=
-# BSD make does it like this.
-echo '.include "confinc.mk" # ignored' > confmf.BSD
-# Other make implementations (GNU, Solaris 10, AIX) do it like this.
-echo 'include confinc.mk # ignored' > confmf.GNU
-_am_result=no
-for s in GNU BSD; do
-  AM_RUN_LOG([${MAKE-make} -f confmf.$s && cat confinc.out])
-  AS_CASE([$?:`cat confinc.out 2>/dev/null`],
-      ['0:this is the am__doit target'],
-      [AS_CASE([$s],
-          [BSD], [am__include='.include' am__quote='"'],
-          [am__include='include' am__quote=''])])
-  if test "$am__include" != "#"; then
-    _am_result="yes ($s style)"
-    break
-  fi
-done
-rm -f confinc.* confmf.*
-AC_MSG_RESULT([${_am_result}])
-AC_SUBST([am__include])])
-AC_SUBST([am__quote])])
+_am_result=none
+# First try GNU make style include.
+echo "include confinc" > confmf
+# Ignore all kinds of additional output from 'make'.
+case `$am_make -s -f confmf 2> /dev/null` in #(
+*the\ am__doit\ target*)
+  am__include=include
+  am__quote=
+  _am_result=GNU
+  ;;
+esac
+# Now try BSD make style include.
+if test "$am__include" = "#"; then
+   echo '.include "confinc"' > confmf
+   case `$am_make -s -f confmf 2> /dev/null` in #(
+   *the\ am__doit\ target*)
+     am__include=.include
+     am__quote="\""
+     _am_result=BSD
+     ;;
+   esac
+fi
+AC_SUBST([am__include])
+AC_SUBST([am__quote])
+AC_MSG_RESULT([$_am_result])
+rm -f confinc confmf
+])
 
 # Fake the existence of programs that GNU maintainers use.  -*- Autoconf -*-
 
-# Copyright (C) 1997-2018 Free Software Foundation, Inc.
+# Copyright (C) 1997-2017 Free Software Foundation, Inc.
 #
 # This file is free software; the Free Software Foundation
 # gives unlimited permission to copy and/or distribute it,
@@ -736,7 +753,7 @@ fi
 
 # Helper functions for option handling.                     -*- Autoconf -*-
 
-# Copyright (C) 2001-2018 Free Software Foundation, Inc.
+# Copyright (C) 2001-2017 Free Software Foundation, Inc.
 #
 # This file is free software; the Free Software Foundation
 # gives unlimited permission to copy and/or distribute it,
@@ -765,7 +782,7 @@ AC_DEFUN([_AM_SET_OPTIONS],
 AC_DEFUN([_AM_IF_OPTION],
 [m4_ifset(_AM_MANGLE_OPTION([$1]), [$2], [$3])])
 
-# Copyright (C) 1999-2018 Free Software Foundation, Inc.
+# Copyright (C) 1999-2017 Free Software Foundation, Inc.
 #
 # This file is free software; the Free Software Foundation
 # gives unlimited permission to copy and/or distribute it,
@@ -812,7 +829,7 @@ AC_LANG_POP([C])])
 # For backward compatibility.
 AC_DEFUN_ONCE([AM_PROG_CC_C_O], [AC_REQUIRE([AC_PROG_CC])])
 
-# Copyright (C) 2001-2018 Free Software Foundation, Inc.
+# Copyright (C) 2001-2017 Free Software Foundation, Inc.
 #
 # This file is free software; the Free Software Foundation
 # gives unlimited permission to copy and/or distribute it,
@@ -831,7 +848,7 @@ AC_DEFUN([AM_RUN_LOG],
 
 # Check to make sure that the build environment is sane.    -*- Autoconf -*-
 
-# Copyright (C) 1996-2018 Free Software Foundation, Inc.
+# Copyright (C) 1996-2017 Free Software Foundation, Inc.
 #
 # This file is free software; the Free Software Foundation
 # gives unlimited permission to copy and/or distribute it,
@@ -912,7 +929,7 @@ AC_CONFIG_COMMANDS_PRE(
 rm -f conftest.file
 ])
 
-# Copyright (C) 2009-2018 Free Software Foundation, Inc.
+# Copyright (C) 2009-2017 Free Software Foundation, Inc.
 #
 # This file is free software; the Free Software Foundation
 # gives unlimited permission to copy and/or distribute it,
@@ -972,7 +989,7 @@ AC_SUBST([AM_BACKSLASH])dnl
 _AM_SUBST_NOTMAKE([AM_BACKSLASH])dnl
 ])
 
-# Copyright (C) 2001-2018 Free Software Foundation, Inc.
+# Copyright (C) 2001-2017 Free Software Foundation, Inc.
 #
 # This file is free software; the Free Software Foundation
 # gives unlimited permission to copy and/or distribute it,
@@ -1000,7 +1017,7 @@ fi
 INSTALL_STRIP_PROGRAM="\$(install_sh) -c -s"
 AC_SUBST([INSTALL_STRIP_PROGRAM])])
 
-# Copyright (C) 2006-2018 Free Software Foundation, Inc.
+# Copyright (C) 2006-2017 Free Software Foundation, Inc.
 #
 # This file is free software; the Free Software Foundation
 # gives unlimited permission to copy and/or distribute it,
@@ -1019,7 +1036,7 @@ AC_DEFUN([AM_SUBST_NOTMAKE], [_AM_SUBST_NOTMAKE($@)])
 
 # Check how to create a tarball.                            -*- Autoconf -*-
 
-# Copyright (C) 2004-2018 Free Software Foundation, Inc.
+# Copyright (C) 2004-2017 Free Software Foundation, Inc.
 #
 # This file is free software; the Free Software Foundation
 # gives unlimited permission to copy and/or distribute it,
diff --git a/libgomp/config/gcn/bar.c b/libgomp/config/gcn/bar.c
index 02fd197..a21529a 100644
--- a/libgomp/config/gcn/bar.c
+++ b/libgomp/config/gcn/bar.c
@@ -43,7 +43,8 @@ gomp_barrier_wait_end (gomp_barrier_t *bar, gomp_barrier_state_t state)
       __atomic_store_n (&bar->generation, bar->generation + BAR_INCR,
 			MEMMODEL_RELAXED);
     }
-  asm ("s_barrier" ::: "memory");
+  if (bar->total > 1)
+    asm ("s_barrier" ::: "memory");
 }
 
 void
@@ -71,7 +72,8 @@ gomp_barrier_wait_last (gomp_barrier_t *bar)
 void
 gomp_team_barrier_wake (gomp_barrier_t *bar, int count)
 {
-  asm ("s_barrier" ::: "memory");
+  if (bar->total > 1)
+    asm ("s_barrier" ::: "memory");
 }
 
 void
@@ -97,7 +99,8 @@ gomp_team_barrier_wait_end (gomp_barrier_t *bar, gomp_barrier_state_t state)
 	  state &= ~BAR_CANCELLED;
 	  state += BAR_INCR - BAR_WAS_LAST;
 	  __atomic_store_n (&bar->generation, state, MEMMODEL_RELAXED);
-	  asm ("s_barrier" ::: "memory");
+	  if (bar->total > 1)
+	    asm ("s_barrier" ::: "memory");
 	  return;
 	}
     }
@@ -172,7 +175,8 @@ gomp_team_barrier_wait_cancel_end (gomp_barrier_t *bar,
 	{
 	  state += BAR_INCR - BAR_WAS_LAST;
 	  __atomic_store_n (&bar->generation, state, MEMMODEL_RELAXED);
-	  asm ("s_barrier" ::: "memory");
+	  if (bar->total > 1)
+	    asm ("s_barrier" ::: "memory");
 	  return false;
 	}
     }
@@ -195,7 +199,8 @@ gomp_team_barrier_wait_cancel_end (gomp_barrier_t *bar,
 	  abort();
 	}
 
-      asm ("s_barrier" ::: "memory");
+      if (bar->total > 1)
+	asm ("s_barrier" ::: "memory");
       gen = __atomic_load_n (&bar->generation, MEMMODEL_RELAXED);
       if (__builtin_expect (gen & BAR_CANCELLED, 0))
 	return true;
diff --git a/libgomp/config/gcn/icv-device.c b/libgomp/config/gcn/icv-device.c
index 1b09f97..ba3a446 100644
--- a/libgomp/config/gcn/icv-device.c
+++ b/libgomp/config/gcn/icv-device.c
@@ -40,6 +40,12 @@ omp_get_default_device (void)
 }
 
 int
+omp_get_initial_device (void)
+{
+  return GOMP_DEVICE_HOST_FALLBACK;
+}
+
+int
 omp_get_num_devices (void)
 {
   return 0;
@@ -66,6 +72,7 @@ omp_is_initial_device (void)
 
 ialias (omp_set_default_device)
 ialias (omp_get_default_device)
+ialias (omp_get_initial_device)
 ialias (omp_get_num_devices)
 ialias (omp_get_num_teams)
 ialias (omp_get_team_num)
diff --git a/libgomp/config/nvptx/bar.c b/libgomp/config/nvptx/bar.c
index 125ca3e..1116561 100644
--- a/libgomp/config/nvptx/bar.c
+++ b/libgomp/config/nvptx/bar.c
@@ -41,7 +41,8 @@ gomp_barrier_wait_end (gomp_barrier_t *bar, gomp_barrier_state_t state)
       __atomic_store_n (&bar->generation, bar->generation + BAR_INCR,
 			MEMMODEL_RELEASE);
     }
-  asm ("bar.sync 1, %0;" : : "r" (32 * bar->total));
+  if (bar->total > 1)
+    asm ("bar.sync 1, %0;" : : "r" (32 * bar->total));
 }
 
 void
@@ -69,7 +70,8 @@ gomp_barrier_wait_last (gomp_barrier_t *bar)
 void
 gomp_team_barrier_wake (gomp_barrier_t *bar, int count)
 {
-  asm ("bar.sync 1, %0;" : : "r" (32 * bar->total));
+  if (bar->total > 1)
+    asm ("bar.sync 1, %0;" : : "r" (32 * bar->total));
 }
 
 void
@@ -95,7 +97,8 @@ gomp_team_barrier_wait_end (gomp_barrier_t *bar, gomp_barrier_state_t state)
 	  state &= ~BAR_CANCELLED;
 	  state += BAR_INCR - BAR_WAS_LAST;
 	  __atomic_store_n (&bar->generation, state, MEMMODEL_RELEASE);
-	  asm ("bar.sync 1, %0;" : : "r" (32 * bar->total));
+	  if (bar->total > 1)
+	    asm ("bar.sync 1, %0;" : : "r" (32 * bar->total));
 	  return;
 	}
     }
@@ -104,7 +107,8 @@ gomp_team_barrier_wait_end (gomp_barrier_t *bar, gomp_barrier_state_t state)
   state &= ~BAR_CANCELLED;
   do
     {
-      asm ("bar.sync 1, %0;" : : "r" (32 * bar->total));
+      if (bar->total > 1)
+	asm ("bar.sync 1, %0;" : : "r" (32 * bar->total));
       gen = __atomic_load_n (&bar->generation, MEMMODEL_ACQUIRE);
       if (__builtin_expect (gen & BAR_TASK_PENDING, 0))
 	{
@@ -158,7 +162,8 @@ gomp_team_barrier_wait_cancel_end (gomp_barrier_t *bar,
 	{
 	  state += BAR_INCR - BAR_WAS_LAST;
 	  __atomic_store_n (&bar->generation, state, MEMMODEL_RELEASE);
-	  asm ("bar.sync 1, %0;" : : "r" (32 * bar->total));
+	  if (bar->total > 1)
+	    asm ("bar.sync 1, %0;" : : "r" (32 * bar->total));
 	  return false;
 	}
     }
@@ -169,7 +174,8 @@ gomp_team_barrier_wait_cancel_end (gomp_barrier_t *bar,
   generation = state;
   do
     {
-      asm ("bar.sync 1, %0;" : : "r" (32 * bar->total));
+      if (bar->total > 1)
+	asm ("bar.sync 1, %0;" : : "r" (32 * bar->total));
       gen = __atomic_load_n (&bar->generation, MEMMODEL_ACQUIRE);
       if (__builtin_expect (gen & BAR_CANCELLED, 0))
 	return true;
diff --git a/libgomp/config/nvptx/icv-device.c b/libgomp/config/nvptx/icv-device.c
index 6a70329..795da7d 100644
--- a/libgomp/config/nvptx/icv-device.c
+++ b/libgomp/config/nvptx/icv-device.c
@@ -40,6 +40,12 @@ omp_get_default_device (void)
 }
 
 int
+omp_get_initial_device (void)
+{
+  return GOMP_DEVICE_HOST_FALLBACK;
+}
+
+int
 omp_get_num_devices (void)
 {
   return 0;
@@ -54,5 +60,6 @@ omp_is_initial_device (void)
 
 ialias (omp_set_default_device)
 ialias (omp_get_default_device)
+ialias (omp_get_initial_device)
 ialias (omp_get_num_devices)
 ialias (omp_is_initial_device)
diff --git a/libgomp/config/t-aix b/libgomp/config/t-aix
index c3bb6c0..f85382f 100644
--- a/libgomp/config/t-aix
+++ b/libgomp/config/t-aix
@@ -1,9 +1,13 @@
 ifeq ($(MULTIBUILDTOP),)
 BITS=$(shell if test -z "`$(CC) -x c -E /dev/null -g3 -o - | grep 64BIT`" ; then \
     echo '64'; else echo '32'; fi)
+ARX=$(shell echo $(AR) | sed -e 's/-X[^ ]*//g')
 MAJOR=$(firstword $(subst :, ,$(libtool_VERSION)))
 all-local:
-	ar -X$(BITS) rc .libs/$(PACKAGE).a ../ppc$(BITS)/$(PACKAGE)/.libs/$(PACKAGE).so.$(MAJOR)
-	ar -X$(BITS) rc ../pthread/$(PACKAGE)/.libs/$(PACKAGE).a ../pthread/ppc$(BITS)/$(PACKAGE)/.libs/$(PACKAGE).so.$(MAJOR)
+	-rm -f .libs/$(PACKAGE).a ../pthread/$(PACKAGE)/.libs/$(PACKAGE).a
+	$(AR) rc .libs/$(PACKAGE).a .libs/$(PACKAGE).so.$(MAJOR)
+	$(AR) rc ../pthread/$(PACKAGE)/.libs/$(PACKAGE).a ../pthread/$(PACKAGE)/.libs/$(PACKAGE).so.$(MAJOR)
+	$(ARX) -X$(BITS) rc .libs/$(PACKAGE).a ../ppc$(BITS)/$(PACKAGE)/.libs/$(PACKAGE).so.$(MAJOR)
+	$(ARX) -X$(BITS) rc ../pthread/$(PACKAGE)/.libs/$(PACKAGE).a ../pthread/ppc$(BITS)/$(PACKAGE)/.libs/$(PACKAGE).so.$(MAJOR)
 endif
 
diff --git a/libgomp/configure b/libgomp/configure
index 534f735..e48371d 100755
--- a/libgomp/configure
+++ b/libgomp/configure
@@ -724,6 +724,7 @@ am__nodep
 AMDEPBACKSLASH
 AMDEP_FALSE
 AMDEP_TRUE
+am__quote
 am__include
 DEPDIR
 OBJEXT
@@ -814,8 +815,7 @@ PACKAGE_VERSION
 PACKAGE_TARNAME
 PACKAGE_NAME
 PATH_SEPARATOR
-SHELL
-am__quote'
+SHELL'
 ac_subst_files=''
 ac_user_opts='
 enable_option_checking
@@ -2885,7 +2885,7 @@ target_alias=${target_alias-$host_alias}
 #  -Wall:  turns on all automake warnings...
 #  -Wno-portability:  ...except this one, since GNU make is required.
 #  -Wno-override: ... and this one, since we do want this in testsuite.
-am__api_version='1.16'
+am__api_version='1.15'
 
 # Find a good install program.  We prefer a C program (faster),
 # so one script is as good as another.  But avoid the broken or
@@ -3401,8 +3401,8 @@ MAKEINFO=${MAKEINFO-"${am_missing_run}makeinfo"}
 
 # For better backward compatibility.  To be removed once Automake 1.9.x
 # dies out for good.  For more background, see:
-# <https://lists.gnu.org/archive/html/automake/2012-07/msg00001.html>
-# <https://lists.gnu.org/archive/html/automake/2012-07/msg00014.html>
+# <http://lists.gnu.org/archive/html/automake/2012-07/msg00001.html>
+# <http://lists.gnu.org/archive/html/automake/2012-07/msg00014.html>
 mkdir_p='$(MKDIR_P)'
 
 # We need awk for the "check" target (and possibly the TAP driver).  The
@@ -3453,7 +3453,7 @@ END
 Aborting the configuration process, to ensure you take notice of the issue.
 
 You can download and install GNU coreutils to get an 'rm' implementation
-that behaves properly: <https://www.gnu.org/software/coreutils/>.
+that behaves properly: <http://www.gnu.org/software/coreutils/>.
 
 If you want to complete the configuration process using your problematic
 'rm' anyway, export the environment variable ACCEPT_INFERIOR_RM_PROGRAM
@@ -4414,45 +4414,45 @@ DEPDIR="${am__leading_dot}deps"
 
 ac_config_commands="$ac_config_commands depfiles"
 
-{ $as_echo "$as_me:${as_lineno-$LINENO}: checking whether ${MAKE-make} supports the include directive" >&5
-$as_echo_n "checking whether ${MAKE-make} supports the include directive... " >&6; }
-cat > confinc.mk << 'END'
+
+am_make=${MAKE-make}
+cat > confinc << 'END'
 am__doit:
-	@echo this is the am__doit target >confinc.out
+	@echo this is the am__doit target
 .PHONY: am__doit
 END
+# If we don't find an include directive, just comment out the code.
+{ $as_echo "$as_me:${as_lineno-$LINENO}: checking for style of include used by $am_make" >&5
+$as_echo_n "checking for style of include used by $am_make... " >&6; }
 am__include="#"
 am__quote=
-# BSD make does it like this.
-echo '.include "confinc.mk" # ignored' > confmf.BSD
-# Other make implementations (GNU, Solaris 10, AIX) do it like this.
-echo 'include confinc.mk # ignored' > confmf.GNU
-_am_result=no
-for s in GNU BSD; do
-  { echo "$as_me:$LINENO: ${MAKE-make} -f confmf.$s && cat confinc.out" >&5
-   (${MAKE-make} -f confmf.$s && cat confinc.out) >&5 2>&5
-   ac_status=$?
-   echo "$as_me:$LINENO: \$? = $ac_status" >&5
-   (exit $ac_status); }
-  case $?:`cat confinc.out 2>/dev/null` in #(
-  '0:this is the am__doit target') :
-    case $s in #(
-  BSD) :
-    am__include='.include' am__quote='"' ;; #(
-  *) :
-    am__include='include' am__quote='' ;;
-esac ;; #(
-  *) :
-     ;;
+_am_result=none
+# First try GNU make style include.
+echo "include confinc" > confmf
+# Ignore all kinds of additional output from 'make'.
+case `$am_make -s -f confmf 2> /dev/null` in #(
+*the\ am__doit\ target*)
+  am__include=include
+  am__quote=
+  _am_result=GNU
+  ;;
 esac
-  if test "$am__include" != "#"; then
-    _am_result="yes ($s style)"
-    break
-  fi
-done
-rm -f confinc.* confmf.*
-{ $as_echo "$as_me:${as_lineno-$LINENO}: result: ${_am_result}" >&5
-$as_echo "${_am_result}" >&6; }
+# Now try BSD make style include.
+if test "$am__include" = "#"; then
+   echo '.include "confinc"' > confmf
+   case `$am_make -s -f confmf 2> /dev/null` in #(
+   *the\ am__doit\ target*)
+     am__include=.include
+     am__quote="\""
+     _am_result=BSD
+     ;;
+   esac
+fi
+
+
+{ $as_echo "$as_me:${as_lineno-$LINENO}: result: $_am_result" >&5
+$as_echo "$_am_result" >&6; }
+rm -f confinc confmf
 
 # Check whether --enable-dependency-tracking was given.
 if test "${enable_dependency_tracking+set}" = set; then :
@@ -17795,7 +17795,7 @@ CC="$CC"
 CXX="$CXX"
 GFORTRAN="$GFORTRAN"
 GDC="$GDC"
-AMDEP_TRUE="$AMDEP_TRUE" MAKE="${MAKE-make}"
+AMDEP_TRUE="$AMDEP_TRUE" ac_aux_dir="$ac_aux_dir"
 
 
 # The HP-UX ksh and POSIX shell print the target directory to stdout
@@ -18785,35 +18785,29 @@ esac ;;
   # Older Autoconf quotes --file arguments for eval, but not when files
   # are listed without --file.  Let's play safe and only enable the eval
   # if we detect the quoting.
-  # TODO: see whether this extra hack can be removed once we start
-  # requiring Autoconf 2.70 or later.
-  case $CONFIG_FILES in #(
-  *\'*) :
-    eval set x "$CONFIG_FILES" ;; #(
-  *) :
-    set x $CONFIG_FILES ;; #(
-  *) :
-     ;;
-esac
+  case $CONFIG_FILES in
+  *\'*) eval set x "$CONFIG_FILES" ;;
+  *)   set x $CONFIG_FILES ;;
+  esac
   shift
-  # Used to flag and report bootstrapping failures.
-  am_rc=0
-  for am_mf
+  for mf
   do
     # Strip MF so we end up with the name of the file.
-    am_mf=`$as_echo "$am_mf" | sed -e 's/:.*$//'`
-    # Check whether this is an Automake generated Makefile which includes
-    # dependency-tracking related rules and includes.
-    # Grep'ing the whole file directly is not great: AIX grep has a line
+    mf=`echo "$mf" | sed -e 's/:.*$//'`
+    # Check whether this is an Automake generated Makefile or not.
+    # We used to match only the files named 'Makefile.in', but
+    # some people rename them; so instead we look at the file content.
+    # Grep'ing the first line is not enough: some people post-process
+    # each Makefile.in and add a new line on top of each file to say so.
+    # Grep'ing the whole file is not good either: AIX grep has a line
     # limit of 2048, but all sed's we know have understand at least 4000.
-    sed -n 's,^am--depfiles:.*,X,p' "$am_mf" | grep X >/dev/null 2>&1 \
-      || continue
-    am_dirpart=`$as_dirname -- "$am_mf" ||
-$as_expr X"$am_mf" : 'X\(.*[^/]\)//*[^/][^/]*/*$' \| \
-	 X"$am_mf" : 'X\(//\)[^/]' \| \
-	 X"$am_mf" : 'X\(//\)$' \| \
-	 X"$am_mf" : 'X\(/\)' \| . 2>/dev/null ||
-$as_echo X"$am_mf" |
+    if sed -n 's,^#.*generated by automake.*,X,p' "$mf" | grep X >/dev/null 2>&1; then
+      dirpart=`$as_dirname -- "$mf" ||
+$as_expr X"$mf" : 'X\(.*[^/]\)//*[^/][^/]*/*$' \| \
+	 X"$mf" : 'X\(//\)[^/]' \| \
+	 X"$mf" : 'X\(//\)$' \| \
+	 X"$mf" : 'X\(/\)' \| . 2>/dev/null ||
+$as_echo X"$mf" |
     sed '/^X\(.*[^/]\)\/\/*[^/][^/]*\/*$/{
 	    s//\1/
 	    q
@@ -18831,48 +18825,53 @@ $as_echo X"$am_mf" |
 	    q
 	  }
 	  s/.*/./; q'`
-    am_filepart=`$as_basename -- "$am_mf" ||
-$as_expr X/"$am_mf" : '.*/\([^/][^/]*\)/*$' \| \
-	 X"$am_mf" : 'X\(//\)$' \| \
-	 X"$am_mf" : 'X\(/\)' \| . 2>/dev/null ||
-$as_echo X/"$am_mf" |
-    sed '/^.*\/\([^/][^/]*\)\/*$/{
+    else
+      continue
+    fi
+    # Extract the definition of DEPDIR, am__include, and am__quote
+    # from the Makefile without running 'make'.
+    DEPDIR=`sed -n 's/^DEPDIR = //p' < "$mf"`
+    test -z "$DEPDIR" && continue
+    am__include=`sed -n 's/^am__include = //p' < "$mf"`
+    test -z "$am__include" && continue
+    am__quote=`sed -n 's/^am__quote = //p' < "$mf"`
+    # Find all dependency output files, they are included files with
+    # $(DEPDIR) in their names.  We invoke sed twice because it is the
+    # simplest approach to changing $(DEPDIR) to its actual value in the
+    # expansion.
+    for file in `sed -n "
+      s/^$am__include $am__quote\(.*(DEPDIR).*\)$am__quote"'$/\1/p' <"$mf" | \
+	 sed -e 's/\$(DEPDIR)/'"$DEPDIR"'/g'`; do
+      # Make sure the directory exists.
+      test -f "$dirpart/$file" && continue
+      fdir=`$as_dirname -- "$file" ||
+$as_expr X"$file" : 'X\(.*[^/]\)//*[^/][^/]*/*$' \| \
+	 X"$file" : 'X\(//\)[^/]' \| \
+	 X"$file" : 'X\(//\)$' \| \
+	 X"$file" : 'X\(/\)' \| . 2>/dev/null ||
+$as_echo X"$file" |
+    sed '/^X\(.*[^/]\)\/\/*[^/][^/]*\/*$/{
 	    s//\1/
 	    q
 	  }
-	  /^X\/\(\/\/\)$/{
+	  /^X\(\/\/\)[^/].*/{
 	    s//\1/
 	    q
 	  }
-	  /^X\/\(\/\).*/{
+	  /^X\(\/\/\)$/{
+	    s//\1/
+	    q
+	  }
+	  /^X\(\/\).*/{
 	    s//\1/
 	    q
 	  }
 	  s/.*/./; q'`
-    { echo "$as_me:$LINENO: cd "$am_dirpart" \
-      && sed -e '/# am--include-marker/d' "$am_filepart" \
-        | $MAKE -f - am--depfiles" >&5
-   (cd "$am_dirpart" \
-      && sed -e '/# am--include-marker/d' "$am_filepart" \
-        | $MAKE -f - am--depfiles) >&5 2>&5
-   ac_status=$?
-   echo "$as_me:$LINENO: \$? = $ac_status" >&5
-   (exit $ac_status); } || am_rc=$?
+      as_dir=$dirpart/$fdir; as_fn_mkdir_p
+      # echo "creating $dirpart/$file"
+      echo '# dummy' > "$dirpart/$file"
+    done
   done
-  if test $am_rc -ne 0; then
-    { { $as_echo "$as_me:${as_lineno-$LINENO}: error: in \`$ac_pwd':" >&5
-$as_echo "$as_me: error: in \`$ac_pwd':" >&2;}
-as_fn_error $? "Something went wrong bootstrapping makefile fragments
-    for automatic dependency tracking.  Try re-running configure with the
-    '--disable-dependency-tracking' option to at least be able to build
-    the package (albeit without support for automatic dependency tracking).
-See \`config.log' for more details" "$LINENO" 5; }
-  fi
-  { am_dirpart=; unset am_dirpart;}
-  { am_filepart=; unset am_filepart;}
-  { am_mf=; unset am_mf;}
-  { am_rc=; unset am_rc;}
-  rm -f conftest-deps.mk
 }
  ;;
     "libtool":C)
diff --git a/libgomp/env.c b/libgomp/env.c
index c0c4730..ab22525 100644
--- a/libgomp/env.c
+++ b/libgomp/env.c
@@ -73,8 +73,10 @@ struct gomp_task_icv gomp_global_icv = {
   .target_data = NULL
 };
 
-unsigned long gomp_max_active_levels_var = INT_MAX;
+unsigned long gomp_max_active_levels_var = gomp_supported_active_levels;
 bool gomp_cancel_var = false;
+enum gomp_target_offload_t gomp_target_offload_var
+  = GOMP_TARGET_OFFLOAD_DEFAULT;
 int gomp_max_task_priority_var = 0;
 #ifndef HAVE_SYNC_BUILTINS
 gomp_mutex_t gomp_managed_threads_lock;
@@ -374,6 +376,44 @@ parse_unsigned_long_list (const char *name, unsigned long *p1stvalue,
   return false;
 }
 
+static void
+parse_target_offload (const char *name, enum gomp_target_offload_t *offload)
+{
+  const char *env;
+  int new_offload = -1;
+
+  env = getenv (name);
+  if (env == NULL)
+    return;
+
+  while (isspace ((unsigned char) *env))
+    ++env;
+  if (strncasecmp (env, "default", 7) == 0)
+    {
+      env += 7;
+      new_offload = GOMP_TARGET_OFFLOAD_DEFAULT;
+    }
+  else if (strncasecmp (env, "mandatory", 9) == 0)
+    {
+      env += 9;
+      new_offload = GOMP_TARGET_OFFLOAD_MANDATORY;
+    }
+  else if (strncasecmp (env, "disabled", 8) == 0)
+    {
+      env += 8;
+      new_offload = GOMP_TARGET_OFFLOAD_DISABLED;
+    }
+  while (isspace ((unsigned char) *env))
+    ++env;
+  if (new_offload != -1 && *env == '\0')
+    {
+      *offload = new_offload;
+      return;
+    }
+
+  gomp_error ("Invalid value for environment variable OMP_TARGET_OFFLOAD");
+}
+
 /* Parse environment variable set to a boolean or list of omp_proc_bind_t
    enum values.  Return true if one was present and it was successfully
    parsed.  */
@@ -1334,6 +1374,21 @@ handle_omp_display_env (unsigned long stacksize, int wait_policy)
     }
   fputs ("'\n", stderr);
 
+  fputs ("  OMP_TARGET_OFFLOAD = '", stderr);
+  switch (gomp_target_offload_var)
+    {
+    case GOMP_TARGET_OFFLOAD_DEFAULT:
+      fputs ("DEFAULT", stderr);
+      break;
+    case GOMP_TARGET_OFFLOAD_MANDATORY:
+      fputs ("MANDATORY", stderr);
+      break;
+    case GOMP_TARGET_OFFLOAD_DISABLED:
+      fputs ("DISABLED", stderr);
+      break;
+    }
+  fputs ("'\n", stderr);
+
   if (verbose)
     {
       fputs ("  GOMP_CPU_AFFINITY = ''\n", stderr);
@@ -1366,9 +1421,12 @@ initialize_env (void)
   parse_boolean ("OMP_CANCELLATION", &gomp_cancel_var);
   parse_boolean ("OMP_DISPLAY_AFFINITY", &gomp_display_affinity_var);
   parse_int ("OMP_DEFAULT_DEVICE", &gomp_global_icv.default_device_var, true);
+  parse_target_offload ("OMP_TARGET_OFFLOAD", &gomp_target_offload_var);
   parse_int ("OMP_MAX_TASK_PRIORITY", &gomp_max_task_priority_var, true);
   parse_unsigned_long ("OMP_MAX_ACTIVE_LEVELS", &gomp_max_active_levels_var,
 		       true);
+  if (gomp_max_active_levels_var > gomp_supported_active_levels)
+    gomp_max_active_levels_var = gomp_supported_active_levels;
   gomp_def_allocator = parse_allocator ();
   if (parse_unsigned_long ("OMP_THREAD_LIMIT", &thread_limit_var, false))
     {
diff --git a/libgomp/fortran.c b/libgomp/fortran.c
index 9d838b3..029dec1 100644
--- a/libgomp/fortran.c
+++ b/libgomp/fortran.c
@@ -63,6 +63,7 @@ ialias_redirect (omp_get_schedule)
 ialias_redirect (omp_get_thread_limit)
 ialias_redirect (omp_set_max_active_levels)
 ialias_redirect (omp_get_max_active_levels)
+ialias_redirect (omp_get_supported_active_levels)
 ialias_redirect (omp_get_level)
 ialias_redirect (omp_get_ancestor_thread_num)
 ialias_redirect (omp_get_team_size)
@@ -418,6 +419,12 @@ omp_get_max_active_levels_ (void)
 }
 
 int32_t
+omp_get_supported_active_levels_ (void)
+{
+  return omp_get_supported_active_levels ();
+}
+
+int32_t
 omp_get_level_ (void)
 {
   return omp_get_level ();
diff --git a/libgomp/icv-device.c b/libgomp/icv-device.c
index 414960c..af7972e 100644
--- a/libgomp/icv-device.c
+++ b/libgomp/icv-device.c
@@ -43,6 +43,12 @@ omp_get_default_device (void)
 }
 
 int
+omp_get_initial_device (void)
+{
+  return gomp_get_num_devices ();
+}
+
+int
 omp_get_num_devices (void)
 {
   return gomp_get_num_devices ();
@@ -57,5 +63,6 @@ omp_is_initial_device (void)
 
 ialias (omp_set_default_device)
 ialias (omp_get_default_device)
+ialias (omp_get_initial_device)
 ialias (omp_get_num_devices)
 ialias (omp_is_initial_device)
diff --git a/libgomp/icv.c b/libgomp/icv.c
index 3c16abb..4da6527 100644
--- a/libgomp/icv.c
+++ b/libgomp/icv.c
@@ -116,7 +116,12 @@ void
 omp_set_max_active_levels (int max_levels)
 {
   if (max_levels >= 0)
-    gomp_max_active_levels_var = max_levels;
+    {
+      if (max_levels <= gomp_supported_active_levels)
+	gomp_max_active_levels_var = max_levels;
+      else
+	gomp_max_active_levels_var = gomp_supported_active_levels;
+    }
 }
 
 int
@@ -126,6 +131,12 @@ omp_get_max_active_levels (void)
 }
 
 int
+omp_get_supported_active_levels (void)
+{
+  return gomp_supported_active_levels;
+}
+
+int
 omp_get_cancellation (void)
 {
   return gomp_cancel_var;
@@ -145,12 +156,6 @@ omp_get_proc_bind (void)
 }
 
 int
-omp_get_initial_device (void)
-{
-  return GOMP_DEVICE_HOST_FALLBACK;
-}
-
-int
 omp_get_num_places (void)
 {
   return gomp_places_list_len;
@@ -227,9 +232,9 @@ ialias (omp_get_max_threads)
 ialias (omp_get_thread_limit)
 ialias (omp_set_max_active_levels)
 ialias (omp_get_max_active_levels)
+ialias (omp_get_supported_active_levels)
 ialias (omp_get_cancellation)
 ialias (omp_get_proc_bind)
-ialias (omp_get_initial_device)
 ialias (omp_get_max_task_priority)
 ialias (omp_get_num_places)
 ialias (omp_get_place_num)
diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h
index 87f939a..da7ac03 100644
--- a/libgomp/libgomp.h
+++ b/libgomp/libgomp.h
@@ -434,12 +434,22 @@ struct gomp_task_icv
   struct target_mem_desc *target_data;
 };
 
+enum gomp_target_offload_t
+{
+  GOMP_TARGET_OFFLOAD_DEFAULT,
+  GOMP_TARGET_OFFLOAD_MANDATORY,
+  GOMP_TARGET_OFFLOAD_DISABLED
+};
+
+#define gomp_supported_active_levels INT_MAX
+
 extern struct gomp_task_icv gomp_global_icv;
 #ifndef HAVE_SYNC_BUILTINS
 extern gomp_mutex_t gomp_managed_threads_lock;
 #endif
 extern unsigned long gomp_max_active_levels_var;
 extern bool gomp_cancel_var;
+extern enum gomp_target_offload_t gomp_target_offload_var;
 extern int gomp_max_task_priority_var;
 extern unsigned long long gomp_spin_count_var, gomp_throttled_spin_count_var;
 extern unsigned long gomp_available_cpus, gomp_managed_threads;
diff --git a/libgomp/libgomp.map b/libgomp/libgomp.map
index c808e810..c5f52f7 100644
--- a/libgomp/libgomp.map
+++ b/libgomp/libgomp.map
@@ -193,6 +193,8 @@ OMP_5.0.1 {
 	omp_destroy_allocator_;
 	omp_alloc;
 	omp_free;
+	omp_get_supported_active_levels;
+	omp_get_supported_active_levels_;
 } OMP_5.0;
 
 GOMP_1.0 {
diff --git a/libgomp/libgomp.texi b/libgomp/libgomp.texi
index 5331230..6937063 100644
--- a/libgomp/libgomp.texi
+++ b/libgomp/libgomp.texi
@@ -166,8 +166,9 @@ linkage, and do not throw exceptions.
 * omp_get_cancellation::        Whether cancellation support is enabled
 * omp_get_default_device::      Get the default device for target regions
 * omp_get_dynamic::             Dynamic teams setting
+* omp_get_initial_device::      Device number of host device
 * omp_get_level::               Number of parallel regions
-* omp_get_max_active_levels::   Maximum number of active regions
+* omp_get_max_active_levels::   Current maximum number of active regions
 * omp_get_max_task_priority::   Maximum task priority value that can be set
 * omp_get_max_threads::         Maximum number of threads of parallel region
 * omp_get_nested::              Nested parallel regions
@@ -177,6 +178,7 @@ linkage, and do not throw exceptions.
 * omp_get_num_threads::         Size of the active team
 * omp_get_proc_bind::           Whether theads may be moved between CPUs
 * omp_get_schedule::            Obtain the runtime scheduling method
+* omp_get_supported_active_levels:: Maximum number of active regions supported
 * omp_get_team_num::            Get team number
 * omp_get_team_size::           Number of threads in a team
 * omp_get_thread_limit::        Maximum number of threads
@@ -352,6 +354,33 @@ disabled by default.
 
 
 
+@node omp_get_initial_device
+@section @code{omp_get_initial_device} -- Return device number of initial device
+@table @asis
+@item @emph{Description}:
+This function returns a device number that represents the host device.
+For OpenMP 5.1, this must be equal to the value returned by the
+@code{omp_get_num_devices} function.
+
+@item @emph{C/C++}
+@multitable @columnfractions .20 .80
+@item @emph{Prototype}: @tab @code{int omp_get_initial_device(void);}
+@end multitable
+
+@item @emph{Fortran}:
+@multitable @columnfractions .20 .80
+@item @emph{Interface}: @tab @code{integer function omp_get_initial_device()}
+@end multitable
+
+@item @emph{See also}:
+@ref{omp_get_num_devices}
+
+@item @emph{Reference}:
+@uref{https://www.openmp.org, OpenMP specification v4.5}, Section 3.2.35.
+@end table
+
+
+
 @node omp_get_level
 @section @code{omp_get_level} -- Obtain the current nesting level
 @table @asis
@@ -379,7 +408,7 @@ which enclose the calling call.
 
 
 @node omp_get_max_active_levels
-@section @code{omp_get_max_active_levels} -- Maximum number of active regions
+@section @code{omp_get_max_active_levels} -- Current maximum number of active regions
 @table @asis
 @item @emph{Description}:
 This function obtains the maximum allowed number of nested, active parallel regions.
@@ -638,6 +667,31 @@ set to the value @code{omp_sched_static}, @code{omp_sched_dynamic},
 @end table
 
 
+@node omp_get_supported_active_levels
+@section @code{omp_get_supported_active_levels} -- Maximum number of active regions supported
+@table @asis
+@item @emph{Description}:
+This function returns the maximum number of nested, active parallel regions
+supported by this implementation.
+
+@item @emph{C/C++}
+@multitable @columnfractions .20 .80
+@item @emph{Prototype}: @tab @code{int omp_get_supported_active_levels(void);}
+@end multitable
+
+@item @emph{Fortran}:
+@multitable @columnfractions .20 .80
+@item @emph{Interface}: @tab @code{integer function omp_get_supported_active_levels()}
+@end multitable
+
+@item @emph{See also}:
+@ref{omp_get_max_active_levels}, @ref{omp_set_max_active_levels}
+
+@item @emph{Reference}:
+@uref{https://www.openmp.org, OpenMP specification v5.0}, Section 3.2.15.
+@end table
+
+
 
 @node omp_get_team_num
 @section @code{omp_get_team_num} -- Get team number
@@ -877,7 +931,8 @@ adjustment of team sizes and @code{false} disables it.
 @table @asis
 @item @emph{Description}:
 This function limits the maximum allowed number of nested, active
-parallel regions.
+parallel regions.  @var{max_levels} must be less or equal to
+the value returned by @code{omp_get_supported_active_levels}.
 
 @item @emph{C/C++}
 @multitable @columnfractions .20 .80
@@ -891,7 +946,8 @@ parallel regions.
 @end multitable
 
 @item @emph{See also}:
-@ref{omp_get_max_active_levels}, @ref{omp_get_active_level}
+@ref{omp_get_max_active_levels}, @ref{omp_get_active_level},
+@ref{omp_get_supported_active_levels}
 
 @item @emph{Reference}:
 @uref{https://www.openmp.org, OpenMP specification v4.5}, Section 3.2.15.
@@ -1353,6 +1409,7 @@ beginning with @env{GOMP_} are GNU extensions.
 * OMP_PLACES::              Specifies on which CPUs the theads should be placed
 * OMP_STACKSIZE::           Set default thread stack size
 * OMP_SCHEDULE::            How threads are scheduled
+* OMP_TARGET_OFFLOAD::      Controls offloading behaviour
 * OMP_THREAD_LIMIT::        Set the maximum number of threads
 * OMP_WAIT_POLICY::         How waiting threads are handled
 * GOMP_CPU_AFFINITY::       Bind threads to specific CPUs
@@ -1626,6 +1683,30 @@ dynamic scheduling and a chunk size of 1 is used.
 
 
 
+@node OMP_TARGET_OFFLOAD
+@section @env{OMP_TARGET_OFFLOAD} -- Controls offloading behaviour
+@cindex Environment Variable
+@cindex Implementation specific setting
+@table @asis
+@item @emph{Description}:
+Specifies the behaviour with regard to offloading code to a device.  This
+variable can be set to one of three values - @code{MANDATORY}, @code{DISABLED}
+or @code{DEFAULT}.
+
+If set to @code{MANDATORY}, the program will terminate with an error if
+the offload device is not present or is not supported.  If set to
+@code{DISABLED}, then offloading is disabled and all code will run on the
+host. If set to @code{DEFAULT}, the program will try offloading to the
+device first, then fall back to running code on the host if it cannot.
+
+If undefined, then the program will behave as if @code{DEFAULT} was set.
+
+@item @emph{Reference}:
+@uref{https://www.openmp.org, OpenMP specification v5.0}, Section 6.17
+@end table
+
+
+
 @node OMP_THREAD_LIMIT
 @section @env{OMP_THREAD_LIMIT} -- Set the maximum number of threads
 @cindex Environment Variable
diff --git a/libgomp/omp.h.in b/libgomp/omp.h.in
index 57af737..a9e6c44 100644
--- a/libgomp/omp.h.in
+++ b/libgomp/omp.h.in
@@ -211,6 +211,7 @@ extern void omp_get_schedule (omp_sched_t *, int *) __GOMP_NOTHROW;
 extern int omp_get_thread_limit (void) __GOMP_NOTHROW;
 extern void omp_set_max_active_levels (int) __GOMP_NOTHROW;
 extern int omp_get_max_active_levels (void) __GOMP_NOTHROW;
+extern int omp_get_supported_active_levels (void) __GOMP_NOTHROW;
 extern int omp_get_level (void) __GOMP_NOTHROW;
 extern int omp_get_ancestor_thread_num (int) __GOMP_NOTHROW;
 extern int omp_get_team_size (int) __GOMP_NOTHROW;
diff --git a/libgomp/omp_lib.f90.in b/libgomp/omp_lib.f90.in
index 3ec31ac..2fae57b 100644
--- a/libgomp/omp_lib.f90.in
+++ b/libgomp/omp_lib.f90.in
@@ -394,6 +394,12 @@
         end interface
 
         interface
+          function omp_get_supported_active_levels ()
+            integer (4) :: omp_get_supported_active_levels
+          end function omp_get_supported_active_levels
+        end interface
+
+        interface
           function omp_get_level ()
             integer (4) :: omp_get_level
           end function omp_get_level
diff --git a/libgomp/omp_lib.h.in b/libgomp/omp_lib.h.in
index 26dbe03..eb1dcc4 100644
--- a/libgomp/omp_lib.h.in
+++ b/libgomp/omp_lib.h.in
@@ -205,9 +205,11 @@
       external omp_get_max_active_levels, omp_get_level
       external omp_get_ancestor_thread_num, omp_get_team_size
       external omp_get_active_level
+      external omp_get_supported_active_levels
       integer(4) omp_get_thread_limit, omp_get_max_active_levels
       integer(4) omp_get_level, omp_get_ancestor_thread_num
       integer(4) omp_get_team_size, omp_get_active_level
+      integer(4) omp_get_supported_active_levels
 
       external omp_in_final
       logical(4) omp_in_final
diff --git a/libgomp/parallel.c b/libgomp/parallel.c
index 2423f11..2fe4f573 100644
--- a/libgomp/parallel.c
+++ b/libgomp/parallel.c
@@ -48,7 +48,14 @@ gomp_resolve_num_threads (unsigned specified, unsigned count)
 
   if (specified == 1)
     return 1;
-  else if (thr->ts.active_level >= 1 && !icv->nest_var)
+
+  if (thr->ts.active_level >= 1
+  /* Accelerators with fixed thread counts require this to return 1 for
+     nested parallel regions.  */
+#if !defined(__AMDGCN__) && !defined(__nvptx__)
+      && !icv->nest_var
+#endif
+      )
     return 1;
   else if (thr->ts.active_level >= gomp_max_active_levels_var)
     return 1;
diff --git a/libgomp/plugin/plugin-nvptx.c b/libgomp/plugin/plugin-nvptx.c
index 390804a..11d4cee 100644
--- a/libgomp/plugin/plugin-nvptx.c
+++ b/libgomp/plugin/plugin-nvptx.c
@@ -701,6 +701,7 @@ link_ptx (CUmodule *module, const struct targ_ptx_obj *ptx_objs,
 
   if (r != CUDA_SUCCESS)
     {
+      GOMP_PLUGIN_error ("Link error log %s\n", &elog[0]);
       GOMP_PLUGIN_error ("cuLinkComplete error: %s", cuda_error (r));
       return false;
     }
@@ -1890,7 +1891,11 @@ nvptx_stacks_free (void *p, int num)
 void
 GOMP_OFFLOAD_run (int ord, void *tgt_fn, void *tgt_vars, void **args)
 {
-  CUfunction function = ((struct targ_fn_descriptor *) tgt_fn)->fn;
+  struct targ_fn_descriptor *tgt_fn_desc
+    = (struct targ_fn_descriptor *) tgt_fn;
+  CUfunction function = tgt_fn_desc->fn;
+  const struct targ_fn_launch *launch = tgt_fn_desc->launch;
+  const char *fn_name = launch->fn;
   CUresult r;
   struct ptx_device *ptx_dev = ptx_devices[ord];
   const char *maybe_abort_msg = "(perhaps abort was called)";
@@ -1925,6 +1930,9 @@ GOMP_OFFLOAD_run (int ord, void *tgt_fn, void *tgt_vars, void **args)
     CU_LAUNCH_PARAM_BUFFER_SIZE, &fn_args_size,
     CU_LAUNCH_PARAM_END
   };
+  GOMP_PLUGIN_debug (0, "  %s: kernel %s: launch"
+		     " [(teams: %u), 1, 1] [(lanes: 32), (threads: %u), 1]\n",
+		     __FUNCTION__, fn_name, teams, threads);
   r = CUDA_CALL_NOCHECK (cuLaunchKernel, function, teams, 1, 1,
 			 32, threads, 1, 0, NULL, NULL, config);
   if (r != CUDA_SUCCESS)
diff --git a/libgomp/target.c b/libgomp/target.c
index ab7ac9b..1a8c67c 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -116,7 +116,15 @@ resolve_device (int device_id)
     }
 
   if (device_id < 0 || device_id >= gomp_get_num_devices ())
-    return NULL;
+    {
+      if (gomp_target_offload_var == GOMP_TARGET_OFFLOAD_MANDATORY
+	  && device_id != GOMP_DEVICE_HOST_FALLBACK
+	  && device_id != num_devices_openmp)
+	gomp_fatal ("OMP_TARGET_OFFLOAD is set to MANDATORY, "
+		    "but device not found");
+
+      return NULL;
+    }
 
   gomp_mutex_lock (&devices[device_id].lock);
   if (devices[device_id].state == GOMP_DEVICE_UNINITIALIZED)
@@ -124,6 +132,11 @@ resolve_device (int device_id)
   else if (devices[device_id].state == GOMP_DEVICE_FINALIZED)
     {
       gomp_mutex_unlock (&devices[device_id].lock);
+
+      if (gomp_target_offload_var == GOMP_TARGET_OFFLOAD_MANDATORY)
+	gomp_fatal ("OMP_TARGET_OFFLOAD is set to MANDATORY, "
+		    "but device is finalized");
+
       return NULL;
     }
   gomp_mutex_unlock (&devices[device_id].lock);
@@ -1997,9 +2010,16 @@ gomp_unload_device (struct gomp_device_descr *devicep)
 /* Host fallback for GOMP_target{,_ext} routines.  */
 
 static void
-gomp_target_fallback (void (*fn) (void *), void **hostaddrs)
+gomp_target_fallback (void (*fn) (void *), void **hostaddrs,
+		      struct gomp_device_descr *devicep)
 {
   struct gomp_thread old_thr, *thr = gomp_thread ();
+
+  if (gomp_target_offload_var == GOMP_TARGET_OFFLOAD_MANDATORY
+      && devicep != NULL)
+    gomp_fatal ("OMP_TARGET_OFFLOAD is set to MANDATORY, but device cannot "
+		"be used for offloading");
+
   old_thr = *thr;
   memset (thr, '\0', sizeof (*thr));
   if (gomp_places_list)
@@ -2107,7 +2127,7 @@ GOMP_target (int device, void (*fn) (void *), const void *unused,
       /* All shared memory devices should use the GOMP_target_ext function.  */
       || devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM
       || !(fn_addr = gomp_get_target_fn_addr (devicep, fn)))
-    return gomp_target_fallback (fn, hostaddrs);
+    return gomp_target_fallback (fn, hostaddrs, devicep);
 
   struct target_mem_desc *tgt_vars
     = gomp_map_vars (devicep, mapnum, hostaddrs, NULL, sizes, kinds, false,
@@ -2243,7 +2263,7 @@ GOMP_target_ext (int device, void (*fn) (void *), size_t mapnum,
 				      tgt_align, tgt_size);
 	    }
 	}
-      gomp_target_fallback (fn, hostaddrs);
+      gomp_target_fallback (fn, hostaddrs, devicep);
       return;
     }
 
@@ -2276,9 +2296,15 @@ GOMP_target_ext (int device, void (*fn) (void *), size_t mapnum,
 /* Host fallback for GOMP_target_data{,_ext} routines.  */
 
 static void
-gomp_target_data_fallback (void)
+gomp_target_data_fallback (struct gomp_device_descr *devicep)
 {
   struct gomp_task_icv *icv = gomp_icv (false);
+
+  if (gomp_target_offload_var == GOMP_TARGET_OFFLOAD_MANDATORY
+      && devicep != NULL)
+    gomp_fatal ("OMP_TARGET_OFFLOAD is set to MANDATORY, but device cannot "
+		"be used for offloading");
+
   if (icv->target_data)
     {
       /* Even when doing a host fallback, if there are any active
@@ -2302,7 +2328,7 @@ GOMP_target_data (int device, const void *unused, size_t mapnum,
   if (devicep == NULL
       || !(devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400)
       || (devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM))
-    return gomp_target_data_fallback ();
+    return gomp_target_data_fallback (devicep);
 
   struct target_mem_desc *tgt
     = gomp_map_vars (devicep, mapnum, hostaddrs, NULL, sizes, kinds, false,
@@ -2321,7 +2347,7 @@ GOMP_target_data_ext (int device, size_t mapnum, void **hostaddrs,
   if (devicep == NULL
       || !(devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400)
       || devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
-    return gomp_target_data_fallback ();
+    return gomp_target_data_fallback (devicep);
 
   struct target_mem_desc *tgt
     = gomp_map_vars (devicep, mapnum, hostaddrs, NULL, sizes, kinds, true,
@@ -2617,7 +2643,7 @@ gomp_target_task_fn (void *data)
 	  || (devicep->can_run_func && !devicep->can_run_func (fn_addr)))
 	{
 	  ttask->state = GOMP_TARGET_TASK_FALLBACK;
-	  gomp_target_fallback (ttask->fn, ttask->hostaddrs);
+	  gomp_target_fallback (ttask->fn, ttask->hostaddrs, devicep);
 	  return false;
 	}
 
@@ -2690,7 +2716,7 @@ GOMP_teams (unsigned int num_teams, unsigned int thread_limit)
 void *
 omp_target_alloc (size_t size, int device_num)
 {
-  if (device_num == GOMP_DEVICE_HOST_FALLBACK)
+  if (device_num == gomp_get_num_devices ())
     return malloc (size);
 
   if (device_num < 0)
@@ -2716,7 +2742,7 @@ omp_target_free (void *device_ptr, int device_num)
   if (device_ptr == NULL)
     return;
 
-  if (device_num == GOMP_DEVICE_HOST_FALLBACK)
+  if (device_num == gomp_get_num_devices ())
     {
       free (device_ptr);
       return;
@@ -2747,7 +2773,7 @@ omp_target_is_present (const void *ptr, int device_num)
   if (ptr == NULL)
     return 1;
 
-  if (device_num == GOMP_DEVICE_HOST_FALLBACK)
+  if (device_num == gomp_get_num_devices ())
     return 1;
 
   if (device_num < 0)
@@ -2781,7 +2807,7 @@ omp_target_memcpy (void *dst, const void *src, size_t length,
   struct gomp_device_descr *dst_devicep = NULL, *src_devicep = NULL;
   bool ret;
 
-  if (dst_device_num != GOMP_DEVICE_HOST_FALLBACK)
+  if (dst_device_num != gomp_get_num_devices ())
     {
       if (dst_device_num < 0)
 	return EINVAL;
@@ -2794,7 +2820,7 @@ omp_target_memcpy (void *dst, const void *src, size_t length,
 	  || dst_devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
 	dst_devicep = NULL;
     }
-  if (src_device_num != GOMP_DEVICE_HOST_FALLBACK)
+  if (src_device_num != num_devices_openmp)
     {
       if (src_device_num < 0)
 	return EINVAL;
@@ -2932,7 +2958,7 @@ omp_target_memcpy_rect (void *dst, const void *src, size_t element_size,
   if (!dst && !src)
     return INT_MAX;
 
-  if (dst_device_num != GOMP_DEVICE_HOST_FALLBACK)
+  if (dst_device_num != gomp_get_num_devices ())
     {
       if (dst_device_num < 0)
 	return EINVAL;
@@ -2945,7 +2971,7 @@ omp_target_memcpy_rect (void *dst, const void *src, size_t element_size,
 	  || dst_devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
 	dst_devicep = NULL;
     }
-  if (src_device_num != GOMP_DEVICE_HOST_FALLBACK)
+  if (src_device_num != num_devices_openmp)
     {
       if (src_device_num < 0)
 	return EINVAL;
@@ -2981,7 +3007,7 @@ int
 omp_target_associate_ptr (const void *host_ptr, const void *device_ptr,
 			  size_t size, size_t device_offset, int device_num)
 {
-  if (device_num == GOMP_DEVICE_HOST_FALLBACK)
+  if (device_num == gomp_get_num_devices ())
     return EINVAL;
 
   if (device_num < 0)
@@ -3044,7 +3070,7 @@ omp_target_associate_ptr (const void *host_ptr, const void *device_ptr,
 int
 omp_target_disassociate_ptr (const void *ptr, int device_num)
 {
-  if (device_num == GOMP_DEVICE_HOST_FALLBACK)
+  if (device_num == gomp_get_num_devices ())
     return EINVAL;
 
   if (device_num < 0)
@@ -3087,9 +3113,9 @@ int
 omp_pause_resource (omp_pause_resource_t kind, int device_num)
 {
   (void) kind;
-  if (device_num == GOMP_DEVICE_HOST_FALLBACK)
+  if (device_num == gomp_get_num_devices ())
     return gomp_pause_host ();
-  if (device_num < 0 || device_num >= gomp_get_num_devices ())
+  if (device_num < 0 || device_num >= num_devices_openmp)
     return -1;
   /* Do nothing for target devices for now.  */
   return 0;
@@ -3253,10 +3279,12 @@ gomp_target_init (void)
   const char *suffix = SONAME_SUFFIX (1);
   const char *cur, *next;
   char *plugin_name;
-  int i, new_num_devices;
+  int i, new_num_devs;
+  int num_devs = 0, num_devs_openmp;
+  struct gomp_device_descr *devs = NULL;
 
-  num_devices = 0;
-  devices = NULL;
+  if (gomp_target_offload_var == GOMP_TARGET_OFFLOAD_DISABLED)
+    return;
 
   cur = OFFLOAD_PLUGINS;
   if (*cur)
@@ -3274,7 +3302,7 @@ gomp_target_init (void)
 	plugin_name = (char *) malloc (prefix_len + cur_len + suffix_len + 1);
 	if (!plugin_name)
 	  {
-	    num_devices = 0;
+	    num_devs = 0;
 	    break;
 	  }
 
@@ -3284,16 +3312,16 @@ gomp_target_init (void)
 
 	if (gomp_load_plugin_for_device (&current_device, plugin_name))
 	  {
-	    new_num_devices = current_device.get_num_devices_func ();
-	    if (new_num_devices >= 1)
+	    new_num_devs = current_device.get_num_devices_func ();
+	    if (new_num_devs >= 1)
 	      {
 		/* Augment DEVICES and NUM_DEVICES.  */
 
-		devices = realloc (devices, (num_devices + new_num_devices)
-				   * sizeof (struct gomp_device_descr));
-		if (!devices)
+		devs = realloc (devs, (num_devs + new_num_devs)
+				      * sizeof (struct gomp_device_descr));
+		if (!devs)
 		  {
-		    num_devices = 0;
+		    num_devs = 0;
 		    free (plugin_name);
 		    break;
 		  }
@@ -3303,12 +3331,12 @@ gomp_target_init (void)
 		current_device.type = current_device.get_type_func ();
 		current_device.mem_map.root = NULL;
 		current_device.state = GOMP_DEVICE_UNINITIALIZED;
-		for (i = 0; i < new_num_devices; i++)
+		for (i = 0; i < new_num_devs; i++)
 		  {
 		    current_device.target_id = i;
-		    devices[num_devices] = current_device;
-		    gomp_mutex_init (&devices[num_devices].lock);
-		    num_devices++;
+		    devs[num_devs] = current_device;
+		    gomp_mutex_init (&devs[num_devs].lock);
+		    num_devs++;
 		  }
 	      }
 	  }
@@ -3320,34 +3348,37 @@ gomp_target_init (void)
 
   /* In DEVICES, sort the GOMP_OFFLOAD_CAP_OPENMP_400 ones first, and set
      NUM_DEVICES_OPENMP.  */
-  struct gomp_device_descr *devices_s
-    = malloc (num_devices * sizeof (struct gomp_device_descr));
-  if (!devices_s)
-    {
-      num_devices = 0;
-      free (devices);
-      devices = NULL;
-    }
-  num_devices_openmp = 0;
-  for (i = 0; i < num_devices; i++)
-    if (devices[i].capabilities & GOMP_OFFLOAD_CAP_OPENMP_400)
-      devices_s[num_devices_openmp++] = devices[i];
-  int num_devices_after_openmp = num_devices_openmp;
-  for (i = 0; i < num_devices; i++)
-    if (!(devices[i].capabilities & GOMP_OFFLOAD_CAP_OPENMP_400))
-      devices_s[num_devices_after_openmp++] = devices[i];
-  free (devices);
-  devices = devices_s;
-
-  for (i = 0; i < num_devices; i++)
+  struct gomp_device_descr *devs_s
+    = malloc (num_devs * sizeof (struct gomp_device_descr));
+  if (!devs_s)
+    {
+      num_devs = 0;
+      free (devs);
+      devs = NULL;
+    }
+  num_devs_openmp = 0;
+  for (i = 0; i < num_devs; i++)
+    if (devs[i].capabilities & GOMP_OFFLOAD_CAP_OPENMP_400)
+      devs_s[num_devs_openmp++] = devs[i];
+  int num_devs_after_openmp = num_devs_openmp;
+  for (i = 0; i < num_devs; i++)
+    if (!(devs[i].capabilities & GOMP_OFFLOAD_CAP_OPENMP_400))
+      devs_s[num_devs_after_openmp++] = devs[i];
+  free (devs);
+  devs = devs_s;
+
+  for (i = 0; i < num_devs; i++)
     {
       /* The 'devices' array can be moved (by the realloc call) until we have
 	 found all the plugins, so registering with the OpenACC runtime (which
 	 takes a copy of the pointer argument) must be delayed until now.  */
-      if (devices[i].capabilities & GOMP_OFFLOAD_CAP_OPENACC_200)
-	goacc_register (&devices[i]);
+      if (devs[i].capabilities & GOMP_OFFLOAD_CAP_OPENACC_200)
+	goacc_register (&devs[i]);
     }
 
+  num_devices = num_devs;
+  num_devices_openmp = num_devs_openmp;
+  devices = devs;
   if (atexit (gomp_target_fini) != 0)
     gomp_fatal ("atexit failed");
 }
diff --git a/libgomp/testsuite/Makefile.in b/libgomp/testsuite/Makefile.in
index 26e925b..e48c3f2 100644
--- a/libgomp/testsuite/Makefile.in
+++ b/libgomp/testsuite/Makefile.in
@@ -1,7 +1,7 @@
-# Makefile.in generated by automake 1.16.1 from Makefile.am.
+# Makefile.in generated by automake 1.15.1 from Makefile.am.
 # @configure_input@
 
-# Copyright (C) 1994-2018 Free Software Foundation, Inc.
+# Copyright (C) 1994-2017 Free Software Foundation, Inc.
 
 # This Makefile.in is free software; the Free Software Foundation
 # gives unlimited permission to copy and/or distribute it,
@@ -331,8 +331,8 @@ Makefile: $(srcdir)/Makefile.in $(top_builddir)/config.status
 	  *config.status*) \
 	    cd $(top_builddir) && $(MAKE) $(AM_MAKEFLAGS) am--refresh;; \
 	  *) \
-	    echo ' cd $(top_builddir) && $(SHELL) ./config.status $(subdir)/$@ $(am__maybe_remake_depfiles)'; \
-	    cd $(top_builddir) && $(SHELL) ./config.status $(subdir)/$@ $(am__maybe_remake_depfiles);; \
+	    echo ' cd $(top_builddir) && $(SHELL) ./config.status $(subdir)/$@ $(am__depfiles_maybe)'; \
+	    cd $(top_builddir) && $(SHELL) ./config.status $(subdir)/$@ $(am__depfiles_maybe);; \
 	esac;
 
 $(top_builddir)/config.status: $(top_srcdir)/configure $(CONFIG_STATUS_DEPENDENCIES)
diff --git a/libgomp/testsuite/libgomp.c++/pr96390.C b/libgomp/testsuite/libgomp.c++/pr96390.C
new file mode 100644
index 0000000..8c770ec
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/pr96390.C
@@ -0,0 +1,49 @@
+/* { dg-additional-options "-O0 -fdump-tree-omplower" } */
+/* { dg-xfail-if "PR 97106/PR 97102 - .alias not (yet) supported for nvptx" { offload_target_nvptx } } */
+
+#include <cstdlib>
+#include <type_traits>
+
+template<int Dim> struct V {
+  int version_called;
+
+  template<bool B = (Dim == 0),
+           typename = typename std::enable_if<B>::type>
+  V ()
+  {
+    version_called = 1;
+  }
+
+  template<typename TArg0,
+           typename = typename std::enable_if<(std::is_same<unsigned long,
+                                               typename std::decay<TArg0>::type>::value)>::type>
+  V (TArg0)
+  {
+    version_called = 2;
+  }
+};
+
+template<int Dim> struct S {
+  V<Dim> v;
+};
+
+int
+main ()
+{
+  int version_set[2] = {-1, -1};
+
+#pragma omp target map(from: version_set[0:2])
+  {
+    S<0> s;
+    version_set[0] = s.v.version_called;
+    V<1> v2((unsigned long) 1);
+    version_set[1] = v2.version_called;
+  }
+
+  if (version_set[0] != 1 || version_set[1] != 2)
+    abort ();
+  return 0;
+}
+
+/* "3" for S<0>::S, V<0>::V<>, and V<1>::V<long unsigned int>:  */
+/* { dg-final { scan-tree-dump-times "__attribute__..omp declare target" 3 "omplower" } } */
diff --git a/libgomp/testsuite/libgomp.c-c++-common/nested-parallel-unbalanced.c b/libgomp/testsuite/libgomp.c-c++-common/nested-parallel-unbalanced.c
new file mode 100644
index 0000000..e777271
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/nested-parallel-unbalanced.c
@@ -0,0 +1,31 @@
+/* Ensure that nested parallel regions work even when the number of loop
+   iterations is not divisible by the number of threads.  */
+
+#include <stdlib.h>
+
+int main() {
+  int A[30][40], B[30][40];
+  size_t n = 30;
+
+  for (size_t i = 0; i < 30; ++i)
+    for (size_t j = 0; j < 40; ++j)
+    A[i][j] = 42;
+
+#pragma omp target map(A[0:30][0:40], B[0:30][0:40])
+  {
+#pragma omp parallel for num_threads(8)
+    for (size_t i = 0; i < n; ++i)
+      {
+#pragma omp parallel for
+	for (size_t j = 0; j < n; ++j)
+	  {
+	    B[i][j] = A[i][j];
+	  }
+      }
+  }
+
+for (size_t i = 0; i < n; ++i)
+  for (size_t j = 0; j < n; ++j)
+    if (B[i][j] != 42)
+      abort ();
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/pr96390.c b/libgomp/testsuite/libgomp.c-c++-common/pr96390.c
new file mode 100644
index 0000000..692bd73
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/pr96390.c
@@ -0,0 +1,26 @@
+/* { dg-additional-options "-O0 -fdump-tree-omplower" } */
+/* { dg-xfail-if "PR 97102/PR 97106 - .alias not (yet) supported for nvptx" { offload_target_nvptx } } */
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+int foo () { return 42; }
+int bar () __attribute__((alias ("foo")));
+int baz () __attribute__((alias ("bar")));
+
+#ifdef __cplusplus
+}
+#endif
+
+
+int
+main ()
+{
+  int n;
+  #pragma omp target map(from:n)
+    n = baz ();
+  if (n != 42)
+    __builtin_abort ();
+}
+/* { dg-final { scan-tree-dump-times "__attribute__..omp declare target" 1 "omplower" } } */
diff --git a/libgomp/testsuite/libgomp.c/lib-2.c b/libgomp/testsuite/libgomp.c/lib-2.c
index 3a3b3f6..ea7a719 100644
--- a/libgomp/testsuite/libgomp.c/lib-2.c
+++ b/libgomp/testsuite/libgomp.c/lib-2.c
@@ -20,6 +20,8 @@ main (void)
   omp_set_max_active_levels (6);
   if (omp_get_max_active_levels () != 6)
     abort ();
+  if (omp_get_max_active_levels () > omp_get_supported_active_levels ())
+    abort ();
 
   return 0;
 }
diff --git a/libgomp/testsuite/libgomp.c/loop-25.c b/libgomp/testsuite/libgomp.c/loop-25.c
new file mode 100644
index 0000000..052da71
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/loop-25.c
@@ -0,0 +1,296 @@
+/* { dg-do run } */
+/* { dg-additional-options "-O2 -fopenmp -fdump-tree-vect-details" } */
+/* { dg-additional-options "-mavx" { target avx_runtime } } */
+
+int x, i, j;
+volatile int a, b, c, d, e, f, g, h;
+int k[11][101];
+extern void abort (void);
+
+int
+main ()
+{
+  int niters, err = 0;
+  for (i = 1; i <= 10; i++)
+    for (j = 1; j <= 10 * i; j++)
+      {
+	k[i][j] = 1;
+	asm volatile ("" : : : "memory");
+      }
+  a = 1; b = 11; c = 1; d = 0; e = 1; f = 10; g = 1; h = 1;
+  niters = 0; i = -100; j = -100; x = -100;
+  #pragma omp parallel for simd collapse(2) lastprivate (i, j, x) reduction(+:niters) reduction(|:err)
+  for (i = 1; i <= 10; i++)
+    for (j = 1; j <= 10 * i; j++)
+      {
+	err |= (i < 1);
+	err |= (i > 10);
+	err |= (j < 1);
+	err |= (j > 10 * i);
+	err |= (k[i][j] != 1);
+	k[i][j]++;
+	x = i * 1024 + (j & 1023);
+	niters++;
+      }
+  if (i != 11 || j != 101 || x != 10340 || niters != 550 || err)
+    abort ();
+  niters = 0; i = -100; j = -100; x = -100;
+  #pragma omp parallel for simd collapse(2) lastprivate (i, j, x) reduction(+:niters) reduction(|:err)
+  for (i = a; i < b; i += c)
+    for (j = d * i + e; j < g + i * f; j += h)
+      {
+	err |= (i < 1);
+	err |= (i > 10);
+	err |= (j < 1);
+	err |= (j > 10 * i);
+	err |= (k[i][j] != 2);
+	k[i][j]++;
+	x = i * 1024 + (j & 1023);
+	niters++;
+      }
+  if (i != 11 || j != 101 || x != 10340 || niters != 550 || err)
+    abort ();
+  for (i = 1; i <= 10; i++)
+    for (j = 1; j <= 10 * i; j++)
+      if (k[i][j] == 3)
+	k[i][j] = 0;
+      else
+	abort ();
+  for (i = 0; i < 11; i++)
+    for (j = 0; j < 101; j++)
+      if (k[i][j] != 0)
+	abort ();
+  for (i = 0; i < 10; i++)
+    for (j = 0; j < 10 * i; j++)
+      {
+        k[i][j] = 1;
+	asm volatile ("" : : : "memory");
+      }
+  a = 0; b = 10; c = 1; d = 0; e = 0; f = 10; g = 0; h = 1;
+  niters = 0; i = -100; j = -100; x = -100;
+  #pragma omp parallel for simd collapse(2) lastprivate (i, j, x) reduction(+:niters) reduction(|:err)
+  for (i = 0; i < 10; i++)
+    for (j = 0; j < 10 * i; j++)
+      {
+	err |= (i < 0);
+	err |= (i >= 10);
+	err |= (j < 0);
+	err |= (j >= 10 * i);
+	err |= (k[i][j] != 1);
+	k[i][j]++;
+	x = i * 1024 + (j & 1023);
+	niters++;
+      }
+  if (i != 10 || j != 90 || x != 9305 || niters != 450 || err)
+    abort ();
+  niters = 0; i = -100; j = -100; x = -100;
+  #pragma omp parallel for simd collapse(2) lastprivate (i, j, x) reduction(+:niters) reduction(|:err)
+  for (i = a; i < b; i += c)
+    for (j = d * i + e; j < g + i * f; j += h)
+      {
+	err |= (i < 0);
+	err |= (i >= 10);
+	err |= (j < 0);
+	err |= (j >= 10 * i);
+	err |= (k[i][j] != 2);
+	k[i][j]++;
+	x = i * 1024 + (j & 1023);
+	niters++;
+      }
+  if (i != 10 || j != 90 || x != 9305 || niters != 450 || err)
+    abort ();
+  for (i = 0; i < 10; i++)
+    for (j = 0; j < 10 * i; j++)
+      if (k[i][j] == 3)
+	k[i][j] = 0;
+      else
+	abort ();
+  for (i = 0; i < 11; i++)
+    for (j = 0; j < 101; j++)
+      if (k[i][j] != 0)
+	abort ();
+  for (i = 4; i < 10; i++)
+    for (j = -9 + 2 * i; j < i; j++)
+      {
+        k[i][j + 1] = 1;
+	asm volatile ("" : : : "memory");
+      }
+  a = 4; b = 10; c = 1; d = 2; e = -9; f = 1; g = 0; h = 1;
+  niters = 0; i = -100; j = -100; x = -100;
+  #pragma omp parallel for simd collapse(2) lastprivate (i, j, x) reduction(+:niters) reduction(|:err)
+  for (i = 4; i < 10; i++)
+    for (j = -9 + 2 * i; j < i; j++)
+      {
+	err |= (i < 4);
+	err |= (i >= 10);
+	err |= (j < -9 + 2 * i);
+	err |= (j >= i);
+	err |= (k[i][j + 1] != 1);
+	k[i][j + 1]++;
+	x = i * 1024 + (j & 1023);
+	niters++;
+      }
+  if (/*i != 10 || j != 9 || */x != 8199 || niters != 15 || err)
+    abort ();
+  niters = 0; i = -100; j = -100; x = -100;
+  #pragma omp parallel for simd collapse(2) lastprivate (i, j, x) reduction(+:niters) reduction(|:err)
+  for (i = a; i < b; i += c)
+    for (j = d * i + e; j < g + i * f; j += h)
+      {
+	err |= (i < 4);
+	err |= (i >= 10);
+	err |= (j < -9 + 2 * i);
+	err |= (j >= i);
+	err |= (k[i][j + 1] != 2);
+	k[i][j + 1]++;
+	x = i * 1024 + (j & 1023);
+	niters++;
+      }
+  if (/*i != 10 || j != 9 || */x != 8199 || niters != 15 || err)
+    abort ();
+  for (i = 4; i < 10; i++)
+    for (j = -9 + 2 * i; j < i; j++)
+      if (k[i][j + 1] == 3)
+	k[i][j + 1] = 0;
+      else
+	abort ();
+  for (i = 0; i < 11; i++)
+    for (j = 0; j < 101; j++)
+      if (k[i][j] != 0)
+	abort ();
+  for (i = 1; i < 10; i += 2)
+    for (j = 1; j < i + 1; j++)
+      {
+	k[i][j] = 1;
+	asm volatile ("" : : : "memory");
+      }
+  a = 1; b = 10; c = 2; d = 0; e = 1; f = 1; g = 1; h = 1;
+  niters = 0; i = -100; j = -100; x = -100;
+  #pragma omp parallel for simd collapse(2) lastprivate (i, j, x) reduction(+:niters) reduction(|:err)
+  for (i = 1; i < 10; i += 2)
+    for (j = 1; j < i + 1; j++)
+      {
+	err |= (i < 1);
+	err |= (i >= 10);
+	err |= (j < 1);
+	err |= (j >= i + 1);
+	err |= (k[i][j] != 1);
+	k[i][j]++;
+	x = i * 1024 + (j & 1023);
+	niters++;
+      }
+  if (i != 11 || j != 10 || x != 9225 || niters != 25 || err)
+    abort ();
+  niters = 0; i = -100; j = -100; x = -100;
+  #pragma omp parallel for simd collapse(2) lastprivate (i, j, x) reduction(+:niters) reduction(|:err)
+  for (i = a; i < b; i += c)
+    for (j = d * i + e; j < g + i * f; j += h)
+      {
+	err |= (i < 1);
+	err |= (i >= 10);
+	err |= (j < 1);
+	err |= (j >= i + 1);
+	err |= (k[i][j] != 2);
+	k[i][j]++;
+	x = i * 1024 + (j & 1023);
+	niters++;
+      }
+  if (i != 11 || j != 10 || x != 9225 || niters != 25 || err)
+    abort ();
+  for (i = 1; i < 10; i += 2)
+    for (j = 1; j < i + 1; j++)
+      if (k[i][j] == 3)
+	k[i][j] = 0;
+      else
+	abort ();
+  for (i = 0; i < 11; i++)
+    for (j = 0; j < 101; j++)
+      if (k[i][j] != 0)
+	abort ();
+  for (j = -11; j >= -41; j -= 15)
+    {
+      k[0][-j] = 1;
+      asm volatile ("" : : : "memory");
+    }
+  a = 4; b = 8; c = 12; d = -8; e = -9; f = -3; g = 6; h = 15;
+  niters = 0; i = -100; j = -100; x = -100;
+  #pragma omp parallel for simd collapse(2) lastprivate (i, j, x) reduction(+:niters) reduction(|:err)
+  for (i = 4; i < 8; i += 12)
+    for (j = -8 * i - 9; j < i * -3 + 6; j += 15)
+      {
+	err |= (i != 4);
+	err |= (j < -41);
+	err |= (j > -11);
+	err |= (k[0][-j] != 1);
+	k[0][-j]++;
+	x = i * 1024 + (j & 1023);
+	niters++;
+      }
+  if (i != 16 || j != 4 || x != 5109 || niters != 3 || err)
+    abort ();
+  niters = 0; i = -100; j = -100; x = -100;
+  #pragma omp parallel for simd collapse(2) lastprivate (i, j, x) reduction(+:niters) reduction(|:err)
+  for (i = a; i < b; i += c)
+    for (j = d * i + e; j < g + i * f; j += h)
+      {
+	err |= (i != 4);
+	err |= (j < -41);
+	err |= (j > -11);
+	err |= (k[0][-j] != 2);
+	k[0][-j]++;
+	x = i * 1024 + (j & 1023);
+	niters++;
+      }
+  if (i != 16 || j != 4 || x != 5109 || niters != 3 || err)
+    abort ();
+  for (j = -11; j >= -41; j -= 15)
+    if (k[0][-j] == 3)
+      k[0][-j] = 0;
+    else
+      abort ();
+  for (j = -11; j >= -41; j--)
+    if (k[0][-j] != 0)
+      abort ();
+  for (j = -34; j <= -7; j++)
+    {
+      k[0][-j] = 1;
+      asm volatile ("" : : : "memory");
+    }
+  a = -13; b = 7; c = 12; d = 3; e = 5; f = 0; g = -6; h = 1;
+  niters = 0; i = -100; j = -100; x = -100;
+  #pragma omp parallel for simd collapse(2) lastprivate (i, j, x) reduction(+:niters) reduction(|:err)
+  for (i = -13; i < 7; i += 12)
+    for (j = 3 * i + 5; j < -6; j++)
+      {
+	err |= (i != -13);
+	err |= (j < -34);
+	err |= (j > -7);
+	err |= (k[0][-j] != 1);
+	k[0][-j]++;
+	x = i * 1024 + (j & 1023);
+	niters++;
+      }
+  if (/*i != 11 || j != 2 || */x != -12295 || niters != 28 || err)
+    abort ();
+  niters = 0; i = -100; j = -100; x = -100;
+  #pragma omp parallel for simd collapse(2) lastprivate (i, j, x) reduction(+:niters) reduction(|:err)
+  for (i = a; i < b; i += c)
+    for (j = d * i + e; j < g + i * f; j += h)
+      {
+	err |= (i != -13);
+	err |= (j < -34);
+	err |= (j > -7);
+	err |= (k[0][-j] != 2);
+	k[0][-j]++;
+	x = i * 1024 + (j & 1023);
+	niters++;
+      }
+  if (/*i != 11 || j != 2 || */x != -12295 || niters != 28 || err)
+    abort ();
+  for (j = -34; j <= -7; j++)
+    if (k[0][-j] == 3)
+      k[0][-j] = 0;
+    else
+      abort ();
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/target-40.c b/libgomp/testsuite/libgomp.c/target-40.c
new file mode 100644
index 0000000..138e162
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/target-40.c
@@ -0,0 +1,10 @@
+#include <omp.h>
+#include <stdlib.h>
+
+int
+main ()
+{
+  if (omp_get_initial_device () != omp_get_num_devices ())
+    abort ();
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/target-41.c b/libgomp/testsuite/libgomp.c/target-41.c
new file mode 100644
index 0000000..d8d756f
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/target-41.c
@@ -0,0 +1,33 @@
+/* { dg-set-target-env-var OMP_TARGET_OFFLOAD "mandatory" } */
+
+#include <omp.h>
+#include <stdlib.h>
+
+int v;
+
+void
+foo (void)
+{
+  v++;
+}
+
+#pragma omp declare target to (v, foo)
+
+int
+main ()
+{
+  /* OMP_TARGET_OFFLOAD=mandatory shouldn't fail for host fallback
+     if it is because the program explicitly asked for the host
+     fallback through if(false) or omp_get_initial_device () as
+     the device.  */
+  #pragma omp target if (v)
+  foo ();
+  #pragma omp target device (omp_get_initial_device ())
+  foo ();
+  omp_set_default_device (omp_get_initial_device ());
+  #pragma omp target
+  foo ();
+  if (v != 3)
+    abort ();
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.fortran/declare-target-3.f90 b/libgomp/testsuite/libgomp.fortran/declare-target-3.f90
new file mode 100644
index 0000000..6e5301d
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/declare-target-3.f90
@@ -0,0 +1,45 @@
+! { dg-additional-options "-fdump-tree-omplower" }
+
+module m
+  implicit none (type, external)
+contains
+  subroutine mod_proc(x)
+    integer :: x(2)
+      x = x + 5
+    end subroutine
+end module m
+
+program main
+  use m
+  implicit none (type, external)
+  if (any (foo() /= [48, 49])) stop 1
+contains
+  integer function fourty_two(y)
+    integer :: y
+    fourty_two = y + 42
+  end function
+
+  integer function wrapper (x, y)
+    integer :: x, y(2)
+    call mod_proc(y)
+    wrapper = fourty_two(x) + 1
+  end function
+
+  function foo()
+    integer :: foo(2)
+    integer :: a(2)
+    integer :: b, summed(2)
+    a = [1, 2]
+    b = -1
+    !$omp target map (tofrom: a, b, summed)
+      summed = wrapper (b, a)
+    !$omp end target
+    if (b /= -1) stop 2            ! unchanged
+    if (any (summed /= 42)) stop 3 ! b + 42 + 1 = 42
+    if (any (a /= [6, 7])) stop 4  ! [1, 2] + 5
+    foo = summed + a               ! [48, 49]
+  end function
+end
+
+! 3 times: mod_proc, fourty_two and wrapper:
+! { dg-final { scan-tree-dump-times "__attribute__..omp declare target" 3 "omplower" } }
diff --git a/libgomp/testsuite/libgomp.fortran/lib4.f90 b/libgomp/testsuite/libgomp.fortran/lib4.f90
index d551cde..5259b3b 100644
--- a/libgomp/testsuite/libgomp.fortran/lib4.f90
+++ b/libgomp/testsuite/libgomp.fortran/lib4.f90
@@ -13,4 +13,6 @@ program lib4
   if (omp_get_thread_limit ().lt.0) stop 3
   call omp_set_max_active_levels (6)
   if (omp_get_max_active_levels ().ne.6) stop 4
+  if (omp_get_max_active_levels () &
+      .gt.omp_get_supported_active_levels ()) stop 5
 end program lib4
diff --git a/libgomp/testsuite/libgomp.fortran/pr66199-5.f90 b/libgomp/testsuite/libgomp.fortran/pr66199-5.f90
index 9482f08..2627a81 100644
--- a/libgomp/testsuite/libgomp.fortran/pr66199-5.f90
+++ b/libgomp/testsuite/libgomp.fortran/pr66199-5.f90
@@ -67,5 +67,5 @@ program main
   if (f1 (0, 1024) /= 1024) stop 1
   if (f2 (0, 1024, 17) /= 1024 + (17 + 5 * 1023)) stop 2
   if (f3 (0, 32, 0, 32) /= 64) stop 3
-  if (f4 (0, 32, 0, 32) /= 64) stop 3
+  if (f4 (0, 32, 0, 32) /= 64) stop 4
 end
diff --git a/libgomp/testsuite/libgomp.fortran/pr95654.f90 b/libgomp/testsuite/libgomp.fortran/pr95654.f90
new file mode 100644
index 0000000..2dddd3d
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/pr95654.f90
@@ -0,0 +1,11 @@
+! { dg-do run }
+program main
+  implicit none
+  integer :: d1
+  !$omp target map(from: d1)
+  !$omp teams distribute parallel do simd default(none) lastprivate(d1) num_teams (2) num_threads (1)
+  do d1 = 0, 31
+  end do
+  !$omp end target
+  if (d1 /= 32) stop 3
+end program main
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-vla.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-vla.c
index 0f51bad..7149357 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-vla.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-vla.c
@@ -59,8 +59,3 @@ main ()
 
   return 0;
 }
-
-
-/* { dg-xfail-run-if "TODO PR90861" { *-*-* } { "-DACC_MEM_SHARED=0" } }
-   This might XPASS if the compiler happens to put the two 'A' VLAs at the same
-   address.  */
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/declare-5.f90 b/libgomp/testsuite/libgomp.oacc-fortran/declare-5.f90
index 2fd25d6..ab434f7 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/declare-5.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/declare-5.f90
@@ -1,4 +1,5 @@
 ! { dg-do run }
+! { dg-xfail-run-if "PR92790 - acc declare device_resident - Fortran common blocks not handled" { *-*-* } { "*" } { "-DACC_DEVICE_TYPE_host=1" } }
 
 module vars
   implicit none