2016-02-10 Basile Starynkevitch <basile@starynkevitch.net>

	{{merging with more of GCC 6, using
	svn merge -r225091:227000 ^/trunk }}
 



git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/melt-branch@233280 138bc75d-0d04-0410-961f-82ee72b054a4

64 files changed, 410 insertions, 524 deletions

diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index d8d37cf7329..1d265b62e99 100644
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,3 +1,99 @@
+2015-08-10  Thomas Schwinge  <thomas@codesourcery.com>
+
+	* libgomp-plugin.h (enum offload_target_type): Remove
+	OFFLOAD_TARGET_TYPE_HOST_NONSHM.
+	* openacc.f90 (openacc_kinds): Remove acc_device_host_nonshm.
+	* openacc.h (enum acc_device_t): Likewise.
+	* openacc_lib.h: Likewise.
+	* oacc-init.c (name_of_acc_device_t): Don't handle it.
+	(acc_on_device): Just use __builtin_acc_on_device.
+	* testsuite/libgomp.oacc-c-c++-common/if-1.c: Don't forbid usage
+	of acc_on_device builtin.
+	* plugin/plugin-host.h: Remove file.
+	* plugin/plugin-host.c: Likewise, but salvage some content into...
+	* oacc-host.c: ... this file.
+	* plugin/Makefrag.am: Don't build libgomp-plugin-host_nonshm.la.
+	* plugin/configfrag.ac (offload_targets): Don't add host_nonshm.
+	* Makefile.in: Regenerate.
+	* configure: Likewise.
+	* testsuite/lib/libgomp.exp
+	(check_effective_target_openacc_host_nonshm_selected): Remove.
+	* testsuite/libgomp.oacc-c++/c++.exp: Don't handle
+	ACC_DEVICE_TYPE=host_nonshm.
+	* testsuite/libgomp.oacc-c/c.exp: Likewise.
+	* testsuite/libgomp.oacc-fortran/fortran.exp: Likewise.
+	* testsuite/libgomp.oacc-c-c++-common/acc_on_device-1.c: Likewise.
+	* testsuite/libgomp.oacc-fortran/acc_on_device-1-1.f90: Likewise.
+	* testsuite/libgomp.oacc-fortran/acc_on_device-1-2.f: Likewise.
+	* testsuite/libgomp.oacc-fortran/acc_on_device-1-3.f: Likewise.
+
+2015-08-10  Thomas Schwinge  <thomas@codesourcery.com>
+	    Jakub Jelinek  <jakub@redhat.com>
+
+	* config/nvptx/affinity.c: New file.
+	* config/nvptx/alloc.c: Likewise.
+	* config/nvptx/bar.c: Likewise.
+	* config/nvptx/barrier.c: Likewise.
+	* config/nvptx/critical.c: Likewise.
+	* config/nvptx/env.c: Likewise.
+	* config/nvptx/error.c: Likewise.
+	* config/nvptx/fortran.c: Likewise.
+	* config/nvptx/iter.c: Likewise.
+	* config/nvptx/iter_ull.c: Likewise.
+	* config/nvptx/libgomp-plugin.c: Likewise.
+	* config/nvptx/lock.c: Likewise.
+	* config/nvptx/loop.c: Likewise.
+	* config/nvptx/loop_ull.c: Likewise.
+	* config/nvptx/mutex.c: Likewise.
+	* config/nvptx/oacc-async.c: Likewise.
+	* config/nvptx/oacc-cuda.c: Likewise.
+	* config/nvptx/oacc-host.c: Likewise.
+	* config/nvptx/oacc-init.c: Likewise.
+	* config/nvptx/oacc-mem.c: Likewise.
+	* config/nvptx/oacc-parallel.c: Likewise.
+	* config/nvptx/oacc-plugin.c: Likewise.
+	* config/nvptx/omp-lock.h: Likewise.
+	* config/nvptx/ordered.c: Likewise.
+	* config/nvptx/parallel.c: Likewise.
+	* config/nvptx/proc.c: Likewise.
+	* config/nvptx/ptrlock.c: Likewise.
+	* config/nvptx/sections.c: Likewise.
+	* config/nvptx/sem.c: Likewise.
+	* config/nvptx/single.c: Likewise.
+	* config/nvptx/splay-tree.c: Likewise.
+	* config/nvptx/target.c: Likewise.
+	* config/nvptx/task.c: Likewise.
+	* config/nvptx/team.c: Likewise.
+	* config/nvptx/time.c: Likewise.
+	* config/nvptx/work.c: Likewise.
+	* configure.ac: Don't probe pthreads support for host nvptx*-*-*.
+	* configure: Regenerate.
+	* configure.tgt (config_path): Set to "nvptx" for target
+	nvptx*-*-*.
+
+2015-08-10  Thomas Schwinge  <thomas@codesourcery.com>
+
+	* testsuite/libgomp.oacc-c-c++-common/vector-type-1.c: New file.
+
+2015-08-03  Nathan Sidwell  <nathan@codesourcery.com>
+
+	* plugin/plugin-nvptx.c: Don't include dlfcn.h.
+	(cuda_errlist): Constify.
+	(errmsg):  Move into ...
+	(cuda_error): ... here.  Make smaller.
+	(_XSTR, _STR): Delete.
+	(cuda_synames): Delete.
+	(verify_device_library): Delete.
+	(nvptx_init): Don't call it.
+
+2015-07-28  Tom de Vries  <tom@codesourcery.com>
+
+	* testsuite/libgomp.c/uns-outer-4.c: New test.
+
+2015-07-24  Cesar Philippidis  <cesar@codesourcery.com>
+
+	* testsuite/libgomp.c/pr66714.c: New test.
+
 2015-07-22  Maxim Blumenthal  <maxim.blumenthal@intel.com>
 
 	PR libgomp/66950
diff --git a/libgomp/Makefile.in b/libgomp/Makefile.in
index 9d07e8d4959..79745ce93f2 100644
--- a/libgomp/Makefile.in
+++ b/libgomp/Makefile.in
@@ -146,15 +146,6 @@ am__installdirs = "$(DESTDIR)$(toolexeclibdir)" "$(DESTDIR)$(infodir)" \
 	"$(DESTDIR)$(fincludedir)" "$(DESTDIR)$(libsubincludedir)" \
 	"$(DESTDIR)$(toolexeclibdir)"
 LTLIBRARIES = $(toolexeclib_LTLIBRARIES)
-libgomp_plugin_host_nonshm_la_DEPENDENCIES = libgomp.la
-am_libgomp_plugin_host_nonshm_la_OBJECTS =  \
-	libgomp_plugin_host_nonshm_la-plugin-host.lo
-libgomp_plugin_host_nonshm_la_OBJECTS =  \
-	$(am_libgomp_plugin_host_nonshm_la_OBJECTS)
-libgomp_plugin_host_nonshm_la_LINK = $(LIBTOOL) --tag=CC \
-	$(libgomp_plugin_host_nonshm_la_LIBTOOLFLAGS) $(LIBTOOLFLAGS) \
-	--mode=link $(CCLD) $(AM_CFLAGS) $(CFLAGS) \
-	$(libgomp_plugin_host_nonshm_la_LDFLAGS) $(LDFLAGS) -o $@
 am__DEPENDENCIES_1 =
 @PLUGIN_NVPTX_TRUE@libgomp_plugin_nvptx_la_DEPENDENCIES = libgomp.la \
 @PLUGIN_NVPTX_TRUE@	$(am__DEPENDENCIES_1)
@@ -196,8 +187,7 @@ FCLD = $(FC)
 FCLINK = $(LIBTOOL) --tag=FC $(AM_LIBTOOLFLAGS) $(LIBTOOLFLAGS) \
 	--mode=link $(FCLD) $(AM_FCFLAGS) $(FCFLAGS) $(AM_LDFLAGS) \
 	$(LDFLAGS) -o $@
-SOURCES = $(libgomp_plugin_host_nonshm_la_SOURCES) \
-	$(libgomp_plugin_nvptx_la_SOURCES) $(libgomp_la_SOURCES)
+SOURCES = $(libgomp_plugin_nvptx_la_SOURCES) $(libgomp_la_SOURCES)
 MULTISRCTOP = 
 MULTIBUILDTOP = 
 MULTIDIRS = 
@@ -401,8 +391,7 @@ libsubincludedir = $(libdir)/gcc/$(target_alias)/$(gcc_version)/include
 AM_CPPFLAGS = $(addprefix -I, $(search_path))
 AM_CFLAGS = $(XCFLAGS)
 AM_LDFLAGS = $(XLDFLAGS) $(SECTION_LDFLAGS) $(OPT_LDFLAGS)
-toolexeclib_LTLIBRARIES = libgomp.la $(am__append_1) \
-	libgomp-plugin-host_nonshm.la
+toolexeclib_LTLIBRARIES = libgomp.la $(am__append_1)
 nodist_toolexeclib_HEADERS = libgomp.spec
 
 # -Wc is only a libtool option.
@@ -437,14 +426,6 @@ libgomp_la_SOURCES = alloc.c barrier.c critical.c env.c error.c iter.c \
 @PLUGIN_NVPTX_TRUE@	$(lt_host_flags) $(PLUGIN_NVPTX_LDFLAGS)
 @PLUGIN_NVPTX_TRUE@libgomp_plugin_nvptx_la_LIBADD = libgomp.la $(PLUGIN_NVPTX_LIBS)
 @PLUGIN_NVPTX_TRUE@libgomp_plugin_nvptx_la_LIBTOOLFLAGS = --tag=disable-static
-libgomp_plugin_host_nonshm_version_info = -version-info $(libtool_VERSION)
-libgomp_plugin_host_nonshm_la_SOURCES = plugin/plugin-host.c
-libgomp_plugin_host_nonshm_la_CPPFLAGS = $(AM_CPPFLAGS) -DHOST_NONSHM_PLUGIN
-libgomp_plugin_host_nonshm_la_LDFLAGS = \
-	$(libgomp_plugin_host_nonshm_version_info) $(lt_host_flags)
-
-libgomp_plugin_host_nonshm_la_LIBADD = libgomp.la
-libgomp_plugin_host_nonshm_la_LIBTOOLFLAGS = --tag=disable-static
 nodist_noinst_HEADERS = libgomp_f.h
 nodist_libsubinclude_HEADERS = omp.h openacc.h
 @USE_FORTRAN_TRUE@nodist_finclude_HEADERS = omp_lib.h omp_lib.f90 omp_lib.mod omp_lib_kinds.mod \
@@ -572,8 +553,6 @@ clean-toolexeclibLTLIBRARIES:
 	  echo "rm -f \"$${dir}/so_locations\""; \
 	  rm -f "$${dir}/so_locations"; \
 	done
-libgomp-plugin-host_nonshm.la: $(libgomp_plugin_host_nonshm_la_OBJECTS) $(libgomp_plugin_host_nonshm_la_DEPENDENCIES) $(EXTRA_libgomp_plugin_host_nonshm_la_DEPENDENCIES) 
-	$(libgomp_plugin_host_nonshm_la_LINK) -rpath $(toolexeclibdir) $(libgomp_plugin_host_nonshm_la_OBJECTS) $(libgomp_plugin_host_nonshm_la_LIBADD) $(LIBS)
 libgomp-plugin-nvptx.la: $(libgomp_plugin_nvptx_la_OBJECTS) $(libgomp_plugin_nvptx_la_DEPENDENCIES) $(EXTRA_libgomp_plugin_nvptx_la_DEPENDENCIES) 
 	$(libgomp_plugin_nvptx_la_LINK) $(am_libgomp_plugin_nvptx_la_rpath) $(libgomp_plugin_nvptx_la_OBJECTS) $(libgomp_plugin_nvptx_la_LIBADD) $(LIBS)
 libgomp.la: $(libgomp_la_OBJECTS) $(libgomp_la_DEPENDENCIES) $(EXTRA_libgomp_la_DEPENDENCIES) 
@@ -596,7 +575,6 @@ distclean-compile:
 @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_host_nonshm_la-plugin-host.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@
@@ -644,13 +622,6 @@ distclean-compile:
 @AMDEP_TRUE@@am__fastdepCC_FALSE@	DEPDIR=$(DEPDIR) $(CCDEPMODE) $(depcomp) @AMDEPBACKSLASH@
 @am__fastdepCC_FALSE@	$(LTCOMPILE) -c -o $@ $<
 
-libgomp_plugin_host_nonshm_la-plugin-host.lo: plugin/plugin-host.c
-@am__fastdepCC_TRUE@	$(LIBTOOL)  --tag=CC $(libgomp_plugin_host_nonshm_la_LIBTOOLFLAGS) $(LIBTOOLFLAGS) --mode=compile $(CC) $(DEFS) $(DEFAULT_INCLUDES) $(INCLUDES) $(libgomp_plugin_host_nonshm_la_CPPFLAGS) $(CPPFLAGS) $(AM_CFLAGS) $(CFLAGS) -MT libgomp_plugin_host_nonshm_la-plugin-host.lo -MD -MP -MF $(DEPDIR)/libgomp_plugin_host_nonshm_la-plugin-host.Tpo -c -o libgomp_plugin_host_nonshm_la-plugin-host.lo `test -f 'plugin/plugin-host.c' || echo '$(srcdir)/'`plugin/plugin-host.c
-@am__fastdepCC_TRUE@	$(am__mv) $(DEPDIR)/libgomp_plugin_host_nonshm_la-plugin-host.Tpo $(DEPDIR)/libgomp_plugin_host_nonshm_la-plugin-host.Plo
-@AMDEP_TRUE@@am__fastdepCC_FALSE@	source='plugin/plugin-host.c' object='libgomp_plugin_host_nonshm_la-plugin-host.lo' libtool=yes @AMDEPBACKSLASH@
-@AMDEP_TRUE@@am__fastdepCC_FALSE@	DEPDIR=$(DEPDIR) $(CCDEPMODE) $(depcomp) @AMDEPBACKSLASH@
-@am__fastdepCC_FALSE@	$(LIBTOOL)  --tag=CC $(libgomp_plugin_host_nonshm_la_LIBTOOLFLAGS) $(LIBTOOLFLAGS) --mode=compile $(CC) $(DEFS) $(DEFAULT_INCLUDES) $(INCLUDES) $(libgomp_plugin_host_nonshm_la_CPPFLAGS) $(CPPFLAGS) $(AM_CFLAGS) $(CFLAGS) -c -o libgomp_plugin_host_nonshm_la-plugin-host.lo `test -f 'plugin/plugin-host.c' || echo '$(srcdir)/'`plugin/plugin-host.c
-
 libgomp_plugin_nvptx_la-plugin-nvptx.lo: plugin/plugin-nvptx.c
 @am__fastdepCC_TRUE@	$(LIBTOOL)  --tag=CC $(libgomp_plugin_nvptx_la_LIBTOOLFLAGS) $(LIBTOOLFLAGS) --mode=compile $(CC) $(DEFS) $(DEFAULT_INCLUDES) $(INCLUDES) $(libgomp_plugin_nvptx_la_CPPFLAGS) $(CPPFLAGS) $(AM_CFLAGS) $(CFLAGS) -MT libgomp_plugin_nvptx_la-plugin-nvptx.lo -MD -MP -MF $(DEPDIR)/libgomp_plugin_nvptx_la-plugin-nvptx.Tpo -c -o libgomp_plugin_nvptx_la-plugin-nvptx.lo `test -f 'plugin/plugin-nvptx.c' || echo '$(srcdir)/'`plugin/plugin-nvptx.c
 @am__fastdepCC_TRUE@	$(am__mv) $(DEPDIR)/libgomp_plugin_nvptx_la-plugin-nvptx.Tpo $(DEPDIR)/libgomp_plugin_nvptx_la-plugin-nvptx.Plo
diff --git a/libgomp/config/nvptx/affinity.c b/libgomp/config/nvptx/affinity.c
new file mode 100644
index 00000000000..e69de29bb2d
--- /dev/null
+++ b/libgomp/config/nvptx/affinity.c
diff --git a/libgomp/config/nvptx/alloc.c b/libgomp/config/nvptx/alloc.c
new file mode 100644
index 00000000000..e69de29bb2d
--- /dev/null
+++ b/libgomp/config/nvptx/alloc.c
diff --git a/libgomp/config/nvptx/bar.c b/libgomp/config/nvptx/bar.c
new file mode 100644
index 00000000000..e69de29bb2d
--- /dev/null
+++ b/libgomp/config/nvptx/bar.c
diff --git a/libgomp/config/nvptx/barrier.c b/libgomp/config/nvptx/barrier.c
new file mode 100644
index 00000000000..e69de29bb2d
--- /dev/null
+++ b/libgomp/config/nvptx/barrier.c
diff --git a/libgomp/config/nvptx/critical.c b/libgomp/config/nvptx/critical.c
new file mode 100644
index 00000000000..e69de29bb2d
--- /dev/null
+++ b/libgomp/config/nvptx/critical.c
diff --git a/libgomp/config/nvptx/env.c b/libgomp/config/nvptx/env.c
new file mode 100644
index 00000000000..e69de29bb2d
--- /dev/null
+++ b/libgomp/config/nvptx/env.c
diff --git a/libgomp/config/nvptx/error.c b/libgomp/config/nvptx/error.c
new file mode 100644
index 00000000000..e69de29bb2d
--- /dev/null
+++ b/libgomp/config/nvptx/error.c
diff --git a/libgomp/config/nvptx/fortran.c b/libgomp/config/nvptx/fortran.c
new file mode 100644
index 00000000000..e69de29bb2d
--- /dev/null
+++ b/libgomp/config/nvptx/fortran.c
diff --git a/libgomp/config/nvptx/iter.c b/libgomp/config/nvptx/iter.c
new file mode 100644
index 00000000000..e69de29bb2d
--- /dev/null
+++ b/libgomp/config/nvptx/iter.c
diff --git a/libgomp/config/nvptx/iter_ull.c b/libgomp/config/nvptx/iter_ull.c
new file mode 100644
index 00000000000..e69de29bb2d
--- /dev/null
+++ b/libgomp/config/nvptx/iter_ull.c
diff --git a/libgomp/config/nvptx/libgomp-plugin.c b/libgomp/config/nvptx/libgomp-plugin.c
new file mode 100644
index 00000000000..e69de29bb2d
--- /dev/null
+++ b/libgomp/config/nvptx/libgomp-plugin.c
diff --git a/libgomp/config/nvptx/lock.c b/libgomp/config/nvptx/lock.c
new file mode 100644
index 00000000000..e69de29bb2d
--- /dev/null
+++ b/libgomp/config/nvptx/lock.c
diff --git a/libgomp/config/nvptx/loop.c b/libgomp/config/nvptx/loop.c
new file mode 100644
index 00000000000..e69de29bb2d
--- /dev/null
+++ b/libgomp/config/nvptx/loop.c
diff --git a/libgomp/config/nvptx/loop_ull.c b/libgomp/config/nvptx/loop_ull.c
new file mode 100644
index 00000000000..e69de29bb2d
--- /dev/null
+++ b/libgomp/config/nvptx/loop_ull.c
diff --git a/libgomp/config/nvptx/mutex.c b/libgomp/config/nvptx/mutex.c
new file mode 100644
index 00000000000..e69de29bb2d
--- /dev/null
+++ b/libgomp/config/nvptx/mutex.c
diff --git a/libgomp/config/nvptx/oacc-async.c b/libgomp/config/nvptx/oacc-async.c
new file mode 100644
index 00000000000..e69de29bb2d
--- /dev/null
+++ b/libgomp/config/nvptx/oacc-async.c
diff --git a/libgomp/config/nvptx/oacc-cuda.c b/libgomp/config/nvptx/oacc-cuda.c
new file mode 100644
index 00000000000..e69de29bb2d
--- /dev/null
+++ b/libgomp/config/nvptx/oacc-cuda.c
diff --git a/libgomp/config/nvptx/oacc-host.c b/libgomp/config/nvptx/oacc-host.c
new file mode 100644
index 00000000000..e69de29bb2d
--- /dev/null
+++ b/libgomp/config/nvptx/oacc-host.c
diff --git a/libgomp/config/nvptx/oacc-init.c b/libgomp/config/nvptx/oacc-init.c
new file mode 100644
index 00000000000..e69de29bb2d
--- /dev/null
+++ b/libgomp/config/nvptx/oacc-init.c
diff --git a/libgomp/config/nvptx/oacc-mem.c b/libgomp/config/nvptx/oacc-mem.c
new file mode 100644
index 00000000000..e69de29bb2d
--- /dev/null
+++ b/libgomp/config/nvptx/oacc-mem.c
diff --git a/libgomp/config/nvptx/oacc-parallel.c b/libgomp/config/nvptx/oacc-parallel.c
new file mode 100644
index 00000000000..e69de29bb2d
--- /dev/null
+++ b/libgomp/config/nvptx/oacc-parallel.c
diff --git a/libgomp/config/nvptx/oacc-plugin.c b/libgomp/config/nvptx/oacc-plugin.c
new file mode 100644
index 00000000000..e69de29bb2d
--- /dev/null
+++ b/libgomp/config/nvptx/oacc-plugin.c
diff --git a/libgomp/config/nvptx/omp-lock.h b/libgomp/config/nvptx/omp-lock.h
new file mode 100644
index 00000000000..2ca7c5e1d60
--- /dev/null
+++ b/libgomp/config/nvptx/omp-lock.h
@@ -0,0 +1,12 @@
+/* This header is used during the build process to find the size and 
+   alignment of the public OpenMP locks, so that we can export data
+   structures without polluting the namespace.
+
+   When using the Linux futex primitive, non-recursive locks require
+   one int.  Recursive locks require we identify the owning task
+   and so require in addition one int and a pointer.  */
+
+typedef int omp_lock_t;
+typedef struct { int lock, count; void *owner; } omp_nest_lock_t;
+typedef int omp_lock_25_t;
+typedef struct { int owner, count; } omp_nest_lock_25_t;
diff --git a/libgomp/config/nvptx/ordered.c b/libgomp/config/nvptx/ordered.c
new file mode 100644
index 00000000000..e69de29bb2d
--- /dev/null
+++ b/libgomp/config/nvptx/ordered.c
diff --git a/libgomp/config/nvptx/parallel.c b/libgomp/config/nvptx/parallel.c
new file mode 100644
index 00000000000..e69de29bb2d
--- /dev/null
+++ b/libgomp/config/nvptx/parallel.c
diff --git a/libgomp/config/nvptx/proc.c b/libgomp/config/nvptx/proc.c
new file mode 100644
index 00000000000..e69de29bb2d
--- /dev/null
+++ b/libgomp/config/nvptx/proc.c
diff --git a/libgomp/config/nvptx/ptrlock.c b/libgomp/config/nvptx/ptrlock.c
new file mode 100644
index 00000000000..e69de29bb2d
--- /dev/null
+++ b/libgomp/config/nvptx/ptrlock.c
diff --git a/libgomp/config/nvptx/sections.c b/libgomp/config/nvptx/sections.c
new file mode 100644
index 00000000000..e69de29bb2d
--- /dev/null
+++ b/libgomp/config/nvptx/sections.c
diff --git a/libgomp/config/nvptx/sem.c b/libgomp/config/nvptx/sem.c
new file mode 100644
index 00000000000..e69de29bb2d
--- /dev/null
+++ b/libgomp/config/nvptx/sem.c
diff --git a/libgomp/config/nvptx/single.c b/libgomp/config/nvptx/single.c
new file mode 100644
index 00000000000..e69de29bb2d
--- /dev/null
+++ b/libgomp/config/nvptx/single.c
diff --git a/libgomp/config/nvptx/splay-tree.c b/libgomp/config/nvptx/splay-tree.c
new file mode 100644
index 00000000000..e69de29bb2d
--- /dev/null
+++ b/libgomp/config/nvptx/splay-tree.c
diff --git a/libgomp/config/nvptx/target.c b/libgomp/config/nvptx/target.c
new file mode 100644
index 00000000000..e69de29bb2d
--- /dev/null
+++ b/libgomp/config/nvptx/target.c
diff --git a/libgomp/config/nvptx/task.c b/libgomp/config/nvptx/task.c
new file mode 100644
index 00000000000..e69de29bb2d
--- /dev/null
+++ b/libgomp/config/nvptx/task.c
diff --git a/libgomp/config/nvptx/team.c b/libgomp/config/nvptx/team.c
new file mode 100644
index 00000000000..e69de29bb2d
--- /dev/null
+++ b/libgomp/config/nvptx/team.c
diff --git a/libgomp/config/nvptx/time.c b/libgomp/config/nvptx/time.c
new file mode 100644
index 00000000000..e69de29bb2d
--- /dev/null
+++ b/libgomp/config/nvptx/time.c
diff --git a/libgomp/config/nvptx/work.c b/libgomp/config/nvptx/work.c
new file mode 100644
index 00000000000..e69de29bb2d
--- /dev/null
+++ b/libgomp/config/nvptx/work.c
diff --git a/libgomp/configure b/libgomp/configure
index f1a92ba9de4..c93e8776241 100755
--- a/libgomp/configure
+++ b/libgomp/configure
@@ -15041,6 +15041,9 @@ case "$host" in
   *-*-rtems*)
     # RTEMS supports Pthreads, but the library is not available at GCC build time.
     ;;
+  nvptx*-*-*)
+    # NVPTX does not support Pthreads, has its own code replacement.
+    ;;
   *)
     # Check to see if -pthread or -lpthread is needed.  Prefer the former.
     # In case the pthread.h system header is not found, this test will fail.
@@ -15167,7 +15170,6 @@ if test x"$plugin_support" = xyes; then
 
 $as_echo "#define PLUGIN_SUPPORT 1" >>confdefs.h
 
-  offload_targets=host_nonshm
 elif test "x${enable_offload_targets-no}" != xno; then
   as_fn_error "Can't support offloading without support for plugins" "$LINENO" 5
 fi
diff --git a/libgomp/configure.ac b/libgomp/configure.ac
index 9cf02189665..b1696d05514 100644
--- a/libgomp/configure.ac
+++ b/libgomp/configure.ac
@@ -179,6 +179,9 @@ case "$host" in
   *-*-rtems*)
     # RTEMS supports Pthreads, but the library is not available at GCC build time.
     ;;
+  nvptx*-*-*)
+    # NVPTX does not support Pthreads, has its own code replacement.
+    ;;
   *)
     # Check to see if -pthread or -lpthread is needed.  Prefer the former.
     # In case the pthread.h system header is not found, this test will fail.
diff --git a/libgomp/configure.tgt b/libgomp/configure.tgt
index 2970f6f7b82..8fad977a211 100644
--- a/libgomp/configure.tgt
+++ b/libgomp/configure.tgt
@@ -151,6 +151,10 @@ case "${target}" in
 	XLDFLAGS="${XLDFLAGS} -lpthread"
 	;;
 
+  nvptx*-*-*)
+	config_path="nvptx"
+	;;
+
   *)
 	;;
 
diff --git a/libgomp/libgomp-plugin.h b/libgomp/libgomp-plugin.h
index 1072ae467be..24fbb9455c9 100644
--- a/libgomp/libgomp-plugin.h
+++ b/libgomp/libgomp-plugin.h
@@ -46,7 +46,7 @@ extern "C" {
 enum offload_target_type
 {
   OFFLOAD_TARGET_TYPE_HOST = 2,
-  OFFLOAD_TARGET_TYPE_HOST_NONSHM = 3,
+  /* OFFLOAD_TARGET_TYPE_HOST_NONSHM = 3 removed.  */
   OFFLOAD_TARGET_TYPE_NVIDIA_PTX = 5,
   OFFLOAD_TARGET_TYPE_INTEL_MIC = 6
 };
diff --git a/libgomp/oacc-host.c b/libgomp/oacc-host.c
index 6dcdbf3658e..17a5102042b 100644
--- a/libgomp/oacc-host.c
+++ b/libgomp/oacc-host.c
@@ -26,51 +26,215 @@
    see the files COPYING3 and COPYING.RUNTIME respectively.  If not, see
    <http://www.gnu.org/licenses/>.  */
 
-/* This shares much of the implementation of the plugin-host.c "host_nonshm"
-   plugin.  */
-#include "plugin/plugin-host.c"
+#include "libgomp.h"
+#include "oacc-int.h"
+
+#include <stdbool.h>
+#include <stddef.h>
+#include <stdint.h>
+
+static struct gomp_device_descr host_dispatch;
+
+static const char *
+host_get_name (void)
+{
+  return host_dispatch.name;
+}
+
+static unsigned int
+host_get_caps (void)
+{
+  return host_dispatch.capabilities;
+}
+
+static int
+host_get_type (void)
+{
+  return host_dispatch.type;
+}
+
+static int
+host_get_num_devices (void)
+{
+  return 1;
+}
+
+static void
+host_init_device (int n __attribute__ ((unused)))
+{
+}
+
+static void
+host_fini_device (int n __attribute__ ((unused)))
+{
+}
+
+static int
+host_load_image (int n __attribute__ ((unused)),
+		 const void *t __attribute__ ((unused)),
+		 struct addr_pair **r __attribute__ ((unused)))
+{
+  return 0;
+}
+
+static void
+host_unload_image (int n __attribute__ ((unused)),
+		   const void *t __attribute__ ((unused)))
+{
+}
+
+static void *
+host_alloc (int n __attribute__ ((unused)), size_t s)
+{
+  return gomp_malloc (s);
+}
+
+static void
+host_free (int n __attribute__ ((unused)), void *p)
+{
+  free (p);
+}
+
+static void *
+host_dev2host (int n __attribute__ ((unused)),
+	       void *h __attribute__ ((unused)),
+	       const void *d __attribute__ ((unused)),
+	       size_t s __attribute__ ((unused)))
+{
+  return NULL;
+}
+
+static void *
+host_host2dev (int n __attribute__ ((unused)),
+	       void *d __attribute__ ((unused)),
+	       const void *h __attribute__ ((unused)),
+	       size_t s __attribute__ ((unused)))
+{
+  return NULL;
+}
+
+static void
+host_run (int n __attribute__ ((unused)), void *fn_ptr, void *vars)
+{
+  void (*fn)(void *) = (void (*)(void *)) fn_ptr;
+
+  fn (vars);
+}
+
+static void
+host_openacc_exec (void (*fn) (void *),
+		   size_t mapnum __attribute__ ((unused)),
+		   void **hostaddrs,
+		   void **devaddrs __attribute__ ((unused)),
+		   size_t *sizes __attribute__ ((unused)),
+		   unsigned short *kinds __attribute__ ((unused)),
+		   int num_gangs __attribute__ ((unused)),
+		   int num_workers __attribute__ ((unused)),
+		   int vector_length __attribute__ ((unused)),
+		   int async __attribute__ ((unused)),
+		   void *targ_mem_desc __attribute__ ((unused)))
+{
+  fn (hostaddrs);
+}
+
+static void
+host_openacc_register_async_cleanup (void *targ_mem_desc __attribute__ ((unused)))
+{
+}
+
+static int
+host_openacc_async_test (int async __attribute__ ((unused)))
+{
+  return 1;
+}
+
+static int
+host_openacc_async_test_all (void)
+{
+  return 1;
+}
+
+static void
+host_openacc_async_wait (int async __attribute__ ((unused)))
+{
+}
+
+static void
+host_openacc_async_wait_async (int async1 __attribute__ ((unused)),
+			       int async2 __attribute__ ((unused)))
+{
+}
+
+static void
+host_openacc_async_wait_all (void)
+{
+}
+
+static void
+host_openacc_async_wait_all_async (int async __attribute__ ((unused)))
+{
+}
+
+static void
+host_openacc_async_set_async (int async __attribute__ ((unused)))
+{
+}
+
+static void *
+host_openacc_create_thread_data (int ord __attribute__ ((unused)))
+{
+  return NULL;
+}
+
+static void
+host_openacc_destroy_thread_data (void *tls_data __attribute__ ((unused)))
+{
+}
 
 static struct gomp_device_descr host_dispatch =
   {
     .name = "host",
-    .capabilities = (GOMP_OFFLOAD_CAP_OPENACC_200
+    .capabilities = (GOMP_OFFLOAD_CAP_SHARED_MEM
 		     | GOMP_OFFLOAD_CAP_NATIVE_EXEC
-		     | GOMP_OFFLOAD_CAP_SHARED_MEM),
+		     | GOMP_OFFLOAD_CAP_OPENACC_200),
     .target_id = 0,
     .type = OFFLOAD_TARGET_TYPE_HOST,
 
-    .get_name_func = GOMP_OFFLOAD_get_name,
-    .get_caps_func = GOMP_OFFLOAD_get_caps,
-    .get_type_func = GOMP_OFFLOAD_get_type,
-    .get_num_devices_func = GOMP_OFFLOAD_get_num_devices,
-    .init_device_func = GOMP_OFFLOAD_init_device,
-    .fini_device_func = GOMP_OFFLOAD_fini_device,
-    .load_image_func = GOMP_OFFLOAD_load_image,
-    .unload_image_func = GOMP_OFFLOAD_unload_image,
-    .alloc_func = GOMP_OFFLOAD_alloc,
-    .free_func = GOMP_OFFLOAD_free,
-    .dev2host_func = GOMP_OFFLOAD_dev2host,
-    .host2dev_func = GOMP_OFFLOAD_host2dev,
-    .run_func = GOMP_OFFLOAD_run,
-
+    .get_name_func = host_get_name,
+    .get_caps_func = host_get_caps,
+    .get_type_func = host_get_type,
+    .get_num_devices_func = host_get_num_devices,
+    .init_device_func = host_init_device,
+    .fini_device_func = host_fini_device,
+    .load_image_func = host_load_image,
+    .unload_image_func = host_unload_image,
+    .alloc_func = host_alloc,
+    .free_func = host_free,
+    .dev2host_func = host_dev2host,
+    .host2dev_func = host_host2dev,
+    .run_func = host_run,
+
+    .mem_map = { NULL },
+    /* .lock initilized in goacc_host_init.  */
     .is_initialized = false,
 
     .openacc = {
-      .exec_func = GOMP_OFFLOAD_openacc_parallel,
+      .data_environ = NULL,
+
+      .exec_func = host_openacc_exec,
 
-      .register_async_cleanup_func
-        = GOMP_OFFLOAD_openacc_register_async_cleanup,
+      .register_async_cleanup_func = host_openacc_register_async_cleanup,
 
-      .async_set_async_func = GOMP_OFFLOAD_openacc_async_set_async,
-      .async_test_func = GOMP_OFFLOAD_openacc_async_test,
-      .async_test_all_func = GOMP_OFFLOAD_openacc_async_test_all,
-      .async_wait_func = GOMP_OFFLOAD_openacc_async_wait,
-      .async_wait_async_func = GOMP_OFFLOAD_openacc_async_wait_async,
-      .async_wait_all_func = GOMP_OFFLOAD_openacc_async_wait_all,
-      .async_wait_all_async_func = GOMP_OFFLOAD_openacc_async_wait_all_async,
+      .async_test_func = host_openacc_async_test,
+      .async_test_all_func = host_openacc_async_test_all,
+      .async_wait_func = host_openacc_async_wait,
+      .async_wait_async_func = host_openacc_async_wait_async,
+      .async_wait_all_func = host_openacc_async_wait_all,
+      .async_wait_all_async_func = host_openacc_async_wait_all_async,
+      .async_set_async_func = host_openacc_async_set_async,
 
-      .create_thread_data_func = GOMP_OFFLOAD_openacc_create_thread_data,
-      .destroy_thread_data_func = GOMP_OFFLOAD_openacc_destroy_thread_data,
+      .create_thread_data_func = host_openacc_create_thread_data,
+      .destroy_thread_data_func = host_openacc_destroy_thread_data,
 
       .cuda = {
 	.get_current_device_func = NULL,
@@ -81,9 +245,9 @@ static struct gomp_device_descr host_dispatch =
     }
   };
 
-/* Register this device type.  */
-static __attribute__ ((constructor))
-void goacc_host_init (void)
+/* Initialize and register this device type.  */
+static __attribute__ ((constructor)) void
+goacc_host_init (void)
 {
   gomp_mutex_init (&host_dispatch.lock);
   goacc_register (&host_dispatch);
diff --git a/libgomp/oacc-init.c b/libgomp/oacc-init.c
index 105d9dc45de..c12f8ad7137 100644
--- a/libgomp/oacc-init.c
+++ b/libgomp/oacc-init.c
@@ -29,7 +29,6 @@
 #include "libgomp.h"
 #include "oacc-int.h"
 #include "openacc.h"
-#include "plugin/plugin-host.h"
 #include <assert.h>
 #include <stdlib.h>
 #include <strings.h>
@@ -102,7 +101,6 @@ name_of_acc_device_t (enum acc_device_t type)
     case acc_device_none: return "none";
     case acc_device_default: return "default";
     case acc_device_host: return "host";
-    case acc_device_host_nonshm: return "host_nonshm";
     case acc_device_not_host: return "not_host";
     case acc_device_nvidia: return "nvidia";
     default: gomp_fatal ("unknown device type %u", (unsigned) type);
@@ -625,18 +623,8 @@ ialias (acc_set_device_num)
 int
 acc_on_device (acc_device_t dev)
 {
-  struct goacc_thread *thr = goacc_thread ();
-
-  /* We only want to appear to be the "host_nonshm" plugin from "offloaded"
-     code -- i.e. within a parallel region.  Test a flag set by the
-     openacc_parallel hook of the host_nonshm plugin to determine that.  */
-  if (acc_get_device_type () == acc_device_host_nonshm
-      && thr && thr->target_tls
-      && ((struct nonshm_thread *)thr->target_tls)->nonshm_exec)
-    return dev == acc_device_host_nonshm || dev == acc_device_not_host;
-
-  /* For OpenACC, libgomp is only built for the host, so this is sufficient.  */
-  return dev == acc_device_host || dev == acc_device_none;
+  /* Just rely on the compiler builtin.  */
+  return __builtin_acc_on_device (dev);
 }
 
 ialias (acc_on_device)
diff --git a/libgomp/openacc.f90 b/libgomp/openacc.f90
index 04d80886a98..fbd63c69b7f 100644
--- a/libgomp/openacc.f90
+++ b/libgomp/openacc.f90
@@ -43,7 +43,7 @@ module openacc_kinds
   integer (acc_device_kind), parameter :: acc_device_none = 0
   integer (acc_device_kind), parameter :: acc_device_default = 1
   integer (acc_device_kind), parameter :: acc_device_host = 2
-  integer (acc_device_kind), parameter :: acc_device_host_nonshm = 3
+  ! integer (acc_device_kind), parameter :: acc_device_host_nonshm = 3 removed.
   integer (acc_device_kind), parameter :: acc_device_not_host = 4
   integer (acc_device_kind), parameter :: acc_device_nvidia = 5
 
diff --git a/libgomp/openacc.h b/libgomp/openacc.h
index 44a1526597d..fc353e1502c 100644
--- a/libgomp/openacc.h
+++ b/libgomp/openacc.h
@@ -53,7 +53,7 @@ typedef enum acc_device_t
     acc_device_none = 0,
     acc_device_default = 1,
     acc_device_host = 2,
-    acc_device_host_nonshm = 3,
+    /* acc_device_host_nonshm = 3 removed.  */
     acc_device_not_host = 4,
     acc_device_nvidia = 5,
     _ACC_device_hwm
diff --git a/libgomp/openacc_lib.h b/libgomp/openacc_lib.h
index 28659a1e0b0..e9c503e1419 100644
--- a/libgomp/openacc_lib.h
+++ b/libgomp/openacc_lib.h
@@ -38,7 +38,8 @@
       integer (acc_device_kind), parameter :: acc_device_none = 0
       integer (acc_device_kind), parameter :: acc_device_default = 1
       integer (acc_device_kind), parameter :: acc_device_host = 2
-      integer (acc_device_kind), parameter :: acc_device_host_nonshm = 3
+!     integer (acc_device_kind), parameter :: acc_device_host_nonshm = 3
+!     removed.
       integer (acc_device_kind), parameter :: acc_device_not_host = 4
       integer (acc_device_kind), parameter :: acc_device_nvidia = 5
 
diff --git a/libgomp/plugin/Makefrag.am b/libgomp/plugin/Makefrag.am
index 167485f52c1..745becd52b7 100644
--- a/libgomp/plugin/Makefrag.am
+++ b/libgomp/plugin/Makefrag.am
@@ -38,12 +38,3 @@ libgomp_plugin_nvptx_la_LDFLAGS += $(PLUGIN_NVPTX_LDFLAGS)
 libgomp_plugin_nvptx_la_LIBADD = libgomp.la $(PLUGIN_NVPTX_LIBS)
 libgomp_plugin_nvptx_la_LIBTOOLFLAGS = --tag=disable-static
 endif
-
-libgomp_plugin_host_nonshm_version_info = -version-info $(libtool_VERSION)
-toolexeclib_LTLIBRARIES += libgomp-plugin-host_nonshm.la
-libgomp_plugin_host_nonshm_la_SOURCES = plugin/plugin-host.c
-libgomp_plugin_host_nonshm_la_CPPFLAGS = $(AM_CPPFLAGS) -DHOST_NONSHM_PLUGIN
-libgomp_plugin_host_nonshm_la_LDFLAGS = \
-	$(libgomp_plugin_host_nonshm_version_info) $(lt_host_flags)
-libgomp_plugin_host_nonshm_la_LIBADD = libgomp.la
-libgomp_plugin_host_nonshm_la_LIBTOOLFLAGS = --tag=disable-static
diff --git a/libgomp/plugin/configfrag.ac b/libgomp/plugin/configfrag.ac
index 254c68853cb..8c2a420d793 100644
--- a/libgomp/plugin/configfrag.ac
+++ b/libgomp/plugin/configfrag.ac
@@ -33,7 +33,6 @@ AC_CHECK_LIB(dl, dlsym, , [plugin_support=no])
 if test x"$plugin_support" = xyes; then
   AC_DEFINE(PLUGIN_SUPPORT, 1,
     [Define if all infrastructure, needed for plugins, is supported.])
-  offload_targets=host_nonshm
 elif test "x${enable_offload_targets-no}" != xno; then
   AC_MSG_ERROR([Can't support offloading without support for plugins])
 fi
diff --git a/libgomp/plugin/plugin-host.c b/libgomp/plugin/plugin-host.c
deleted file mode 100644
index da3c5f4cdb7..00000000000
--- a/libgomp/plugin/plugin-host.c
+++ /dev/null
@@ -1,259 +0,0 @@
-/* OpenACC Runtime Library: acc_device_host, acc_device_host_nonshm.
-
-   Copyright (C) 2013-2015 Free Software Foundation, Inc.
-
-   Contributed by Mentor Embedded.
-
-   This file is part of the GNU Offloading and Multi Processing Library
-   (libgomp).
-
-   Libgomp is free software; you can redistribute it and/or modify it
-   under the terms of the GNU General Public License as published by
-   the Free Software Foundation; either version 3, or (at your option)
-   any later version.
-
-   Libgomp is distributed in the hope that it will be useful, but WITHOUT ANY
-   WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS
-   FOR A PARTICULAR PURPOSE.  See the GNU General Public License for
-   more details.
-
-   Under Section 7 of GPL version 3, you are granted additional
-   permissions described in the GCC Runtime Library Exception, version
-   3.1, as published by the Free Software Foundation.
-
-   You should have received a copy of the GNU General Public License and
-   a copy of the GCC Runtime Library Exception along with this program;
-   see the files COPYING3 and COPYING.RUNTIME respectively.  If not, see
-   <http://www.gnu.org/licenses/>.  */
-
-/* Simple implementation of support routines for a shared-memory
-   acc_device_host, and a non-shared memory acc_device_host_nonshm, with the
-   latter built as a plugin.  */
-
-#include "openacc.h"
-#include "config.h"
-#ifdef HOST_NONSHM_PLUGIN
-#include "libgomp-plugin.h"
-#include "oacc-plugin.h"
-#else
-#include "libgomp.h"
-#include "oacc-int.h"
-#endif
-
-#include <stdint.h>
-#include <stdlib.h>
-#include <string.h>
-#include <stdio.h>
-#include <stdbool.h>
-
-#ifdef HOST_NONSHM_PLUGIN
-#define STATIC
-#define GOMP(X) GOMP_PLUGIN_##X
-#define SELF "host_nonshm plugin: "
-#else
-#define STATIC static
-#define GOMP(X) gomp_##X
-#define SELF "host: "
-#endif
-
-#ifdef HOST_NONSHM_PLUGIN
-#include "plugin-host.h"
-#endif
-
-STATIC const char *
-GOMP_OFFLOAD_get_name (void)
-{
-#ifdef HOST_NONSHM_PLUGIN
-  return "host_nonshm";
-#else
-  return "host";
-#endif
-}
-
-STATIC unsigned int
-GOMP_OFFLOAD_get_caps (void)
-{
-  unsigned int caps = (GOMP_OFFLOAD_CAP_OPENACC_200
-		       | GOMP_OFFLOAD_CAP_NATIVE_EXEC);
-
-#ifndef HOST_NONSHM_PLUGIN
-  caps |= GOMP_OFFLOAD_CAP_SHARED_MEM;
-#endif
-
-  return caps;
-}
-
-STATIC int
-GOMP_OFFLOAD_get_type (void)
-{
-#ifdef HOST_NONSHM_PLUGIN
-  return OFFLOAD_TARGET_TYPE_HOST_NONSHM;
-#else
-  return OFFLOAD_TARGET_TYPE_HOST;
-#endif
-}
-
-STATIC int
-GOMP_OFFLOAD_get_num_devices (void)
-{
-  return 1;
-}
-
-STATIC void
-GOMP_OFFLOAD_init_device (int n __attribute__ ((unused)))
-{
-}
-
-STATIC void
-GOMP_OFFLOAD_fini_device (int n __attribute__ ((unused)))
-{
-}
-
-STATIC int
-GOMP_OFFLOAD_load_image (int n __attribute__ ((unused)),
-			 const void *t __attribute__ ((unused)),
-			 struct addr_pair **r __attribute__ ((unused)))
-{
-  return 0;
-}
-
-STATIC void
-GOMP_OFFLOAD_unload_image (int n __attribute__ ((unused)),
-			   const void *t __attribute__ ((unused)))
-{
-}
-
-STATIC void *
-GOMP_OFFLOAD_alloc (int n __attribute__ ((unused)), size_t s)
-{
-  return GOMP (malloc) (s);
-}
-
-STATIC void
-GOMP_OFFLOAD_free (int n __attribute__ ((unused)), void *p)
-{
-  free (p);
-}
-
-STATIC void *
-GOMP_OFFLOAD_host2dev (int n __attribute__ ((unused)), void *d, const void *h,
-		       size_t s)
-{
-#ifdef HOST_NONSHM_PLUGIN
-  memcpy (d, h, s);
-#endif
-
-  return 0;
-}
-
-STATIC void *
-GOMP_OFFLOAD_dev2host (int n __attribute__ ((unused)), void *h, const void *d,
-		       size_t s)
-{
-#ifdef HOST_NONSHM_PLUGIN
-  memcpy (h, d, s);
-#endif
-
-  return 0;
-}
-
-STATIC void
-GOMP_OFFLOAD_run (int n __attribute__ ((unused)), void *fn_ptr, void *vars)
-{
-  void (*fn)(void *) = (void (*)(void *)) fn_ptr;
-
-  fn (vars);
-}
-
-STATIC void
-GOMP_OFFLOAD_openacc_parallel (void (*fn) (void *),
-			       size_t mapnum __attribute__ ((unused)),
-			       void **hostaddrs __attribute__ ((unused)),
-			       void **devaddrs __attribute__ ((unused)),
-			       size_t *sizes __attribute__ ((unused)),
-			       unsigned short *kinds __attribute__ ((unused)),
-			       int num_gangs __attribute__ ((unused)),
-			       int num_workers __attribute__ ((unused)),
-			       int vector_length __attribute__ ((unused)),
-			       int async __attribute__ ((unused)),
-			       void *targ_mem_desc __attribute__ ((unused)))
-{
-#ifdef HOST_NONSHM_PLUGIN
-  struct nonshm_thread *thd = GOMP_PLUGIN_acc_thread ();
-  thd->nonshm_exec = true;
-  fn (devaddrs);
-  thd->nonshm_exec = false;
-#else
-  fn (hostaddrs);
-#endif
-}
-
-STATIC void
-GOMP_OFFLOAD_openacc_register_async_cleanup (void *targ_mem_desc)
-{
-#ifdef HOST_NONSHM_PLUGIN
-  /* "Asynchronous" launches are executed synchronously on the (non-SHM) host,
-     so there's no point in delaying host-side cleanup -- just do it now.  */
-  GOMP_PLUGIN_async_unmap_vars (targ_mem_desc);
-#endif
-}
-
-STATIC void
-GOMP_OFFLOAD_openacc_async_set_async (int async __attribute__ ((unused)))
-{
-}
-
-STATIC int
-GOMP_OFFLOAD_openacc_async_test (int async __attribute__ ((unused)))
-{
-  return 1;
-}
-
-STATIC int
-GOMP_OFFLOAD_openacc_async_test_all (void)
-{
-  return 1;
-}
-
-STATIC void
-GOMP_OFFLOAD_openacc_async_wait (int async __attribute__ ((unused)))
-{
-}
-
-STATIC void
-GOMP_OFFLOAD_openacc_async_wait_all (void)
-{
-}
-
-STATIC void
-GOMP_OFFLOAD_openacc_async_wait_async (int async1 __attribute__ ((unused)),
-				       int async2 __attribute__ ((unused)))
-{
-}
-
-STATIC void
-GOMP_OFFLOAD_openacc_async_wait_all_async (int async __attribute__ ((unused)))
-{
-}
-
-STATIC void *
-GOMP_OFFLOAD_openacc_create_thread_data (int ord
-					 __attribute__ ((unused)))
-{
-#ifdef HOST_NONSHM_PLUGIN
-  struct nonshm_thread *thd
-    = GOMP_PLUGIN_malloc (sizeof (struct nonshm_thread));
-  thd->nonshm_exec = false;
-  return thd;
-#else
-  return NULL;
-#endif
-}
-
-STATIC void
-GOMP_OFFLOAD_openacc_destroy_thread_data (void *tls_data)
-{
-#ifdef HOST_NONSHM_PLUGIN
-  free (tls_data);
-#endif
-}
diff --git a/libgomp/plugin/plugin-host.h b/libgomp/plugin/plugin-host.h
deleted file mode 100644
index 96955d19414..00000000000
--- a/libgomp/plugin/plugin-host.h
+++ /dev/null
@@ -1,37 +0,0 @@
-/* OpenACC Runtime Library: acc_device_host, acc_device_host_nonshm.
-
-   Copyright (C) 2015 Free Software Foundation, Inc.
-
-   Contributed by Mentor Embedded.
-
-   This file is part of the GNU Offloading and Multi Processing Library
-   (libgomp).
-
-   Libgomp is free software; you can redistribute it and/or modify it
-   under the terms of the GNU General Public License as published by
-   the Free Software Foundation; either version 3, or (at your option)
-   any later version.
-
-   Libgomp is distributed in the hope that it will be useful, but WITHOUT ANY
-   WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS
-   FOR A PARTICULAR PURPOSE.  See the GNU General Public License for
-   more details.
-
-   Under Section 7 of GPL version 3, you are granted additional
-   permissions described in the GCC Runtime Library Exception, version
-   3.1, as published by the Free Software Foundation.
-
-   You should have received a copy of the GNU General Public License and
-   a copy of the GCC Runtime Library Exception along with this program;
-   see the files COPYING3 and COPYING.RUNTIME respectively.  If not, see
-   <http://www.gnu.org/licenses/>.  */
-
-#ifndef PLUGIN_HOST_H
-#define PLUGIN_HOST_H
-
-struct nonshm_thread
-{
-  bool nonshm_exec;
-};
-
-#endif
diff --git a/libgomp/plugin/plugin-nvptx.c b/libgomp/plugin/plugin-nvptx.c
index fc296321812..d02a3fd4b9c 100644
--- a/libgomp/plugin/plugin-nvptx.c
+++ b/libgomp/plugin/plugin-nvptx.c
@@ -43,16 +43,15 @@
 #include <stdint.h>
 #include <string.h>
 #include <stdio.h>
-#include <dlfcn.h>
 #include <unistd.h>
 #include <assert.h>
 
 #define	ARRAYSIZE(X) (sizeof (X) / sizeof ((X)[0]))
 
-static struct
+static const struct
 {
   CUresult r;
-  char *m;
+  const char *m;
 } cuda_errlist[]=
 {
   { CUDA_ERROR_INVALID_VALUE, "invalid value" },
@@ -109,9 +108,7 @@ static struct
   { CUDA_ERROR_UNKNOWN, "unknown" }
 };
 
-static char errmsg[128];
-
-static char *
+static const char *
 cuda_error (CUresult r)
 {
   int i;
@@ -119,12 +116,14 @@ cuda_error (CUresult r)
   for (i = 0; i < ARRAYSIZE (cuda_errlist); i++)
     {
       if (cuda_errlist[i].r == r)
-	return &cuda_errlist[i].m[0];
+	return cuda_errlist[i].m;
     }
 
-  sprintf (&errmsg[0], "unknown result code: %5d", r);
+  static char errmsg[30];
+
+  snprintf (errmsg, sizeof (errmsg), "unknown error code: %d", r);
 
-  return &errmsg[0];
+  return errmsg;
 }
 
 static unsigned int instantiated_devices = 0;
@@ -353,74 +352,6 @@ static struct ptx_event *ptx_events;
 
 static struct ptx_device **ptx_devices;
 
-#define _XSTR(s) _STR(s)
-#define _STR(s) #s
-
-static struct _synames
-{
-  char *n;
-} cuda_symnames[] =
-{
-  { _XSTR (cuCtxCreate) },
-  { _XSTR (cuCtxDestroy) },
-  { _XSTR (cuCtxGetCurrent) },
-  { _XSTR (cuCtxPushCurrent) },
-  { _XSTR (cuCtxSynchronize) },
-  { _XSTR (cuDeviceGet) },
-  { _XSTR (cuDeviceGetAttribute) },
-  { _XSTR (cuDeviceGetCount) },
-  { _XSTR (cuEventCreate) },
-  { _XSTR (cuEventDestroy) },
-  { _XSTR (cuEventQuery) },
-  { _XSTR (cuEventRecord) },
-  { _XSTR (cuInit) },
-  { _XSTR (cuLaunchKernel) },
-  { _XSTR (cuLinkAddData) },
-  { _XSTR (cuLinkComplete) },
-  { _XSTR (cuLinkCreate) },
-  { _XSTR (cuMemAlloc) },
-  { _XSTR (cuMemAllocHost) },
-  { _XSTR (cuMemcpy) },
-  { _XSTR (cuMemcpyDtoH) },
-  { _XSTR (cuMemcpyDtoHAsync) },
-  { _XSTR (cuMemcpyHtoD) },
-  { _XSTR (cuMemcpyHtoDAsync) },
-  { _XSTR (cuMemFree) },
-  { _XSTR (cuMemFreeHost) },
-  { _XSTR (cuMemGetAddressRange) },
-  { _XSTR (cuMemHostGetDevicePointer) },
-  { _XSTR (cuMemHostRegister) },
-  { _XSTR (cuMemHostUnregister) },
-  { _XSTR (cuModuleGetFunction) },
-  { _XSTR (cuModuleLoadData) },
-  { _XSTR (cuStreamDestroy) },
-  { _XSTR (cuStreamQuery) },
-  { _XSTR (cuStreamSynchronize) },
-  { _XSTR (cuStreamWaitEvent) }
-};
-
-static int
-verify_device_library (void)
-{
-  int i;
-  void *dh, *ds;
-
-  dh = dlopen ("libcuda.so", RTLD_LAZY);
-  if (!dh)
-    return -1;
-
-  for (i = 0; i < ARRAYSIZE (cuda_symnames); i++)
-    {
-      ds = dlsym (dh, cuda_symnames[i].n);
-      if (!ds)
-        return -1;
-    }
-
-  dlclose (dh);
-
-  return 0;
-}
-
 static inline struct nvptx_thread *
 nvptx_thread (void)
 {
@@ -601,16 +532,11 @@ static bool
 nvptx_init (void)
 {
   CUresult r;
-  int rc;
   int ndevs;
 
   if (instantiated_devices != 0)
     return true;
 
-  rc = verify_device_library ();
-  if (rc < 0)
-    return false;
-
   r = cuInit (0);
   if (r != CUDA_SUCCESS)
     GOMP_PLUGIN_fatal ("cuInit error: %s", cuda_error (r));
diff --git a/libgomp/testsuite/lib/libgomp.exp b/libgomp/testsuite/lib/libgomp.exp
index 438777f1fc4..f04b163fafc 100644
--- a/libgomp/testsuite/lib/libgomp.exp
+++ b/libgomp/testsuite/lib/libgomp.exp
@@ -355,13 +355,3 @@ proc check_effective_target_openacc_host_selected { } {
     }
     return 0;
 }
-
-# Return 1 if the host_nonshm target is selected for offloaded
-
-proc check_effective_target_openacc_host_nonshm_selected { } {
-    global offload_target_openacc
-    if { $offload_target_openacc == "host_nonshm" } {
-        return 1;
-    }
-    return 0;
-}
diff --git a/libgomp/testsuite/libgomp.c/pr66714.c b/libgomp/testsuite/libgomp.c/pr66714.c
new file mode 100644
index 00000000000..c9af4a9b4a2
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/pr66714.c
@@ -0,0 +1,17 @@
+/* { dg-do "compile" } */
+/* { dg-additional-options "--param ggc-min-expand=0" } */
+/* { dg-additional-options "--param ggc-min-heapsize=0" } */
+/* { dg-additional-options "-g" } */
+
+/* Minimized from on target-2.c.  */
+
+void
+fn3 (int x)
+{
+  double b[3 * x];
+  int i;
+#pragma omp target
+#pragma omp parallel for
+  for (i = 0; i < x; i++)
+    b[i] += 1;
+}
diff --git a/libgomp/testsuite/libgomp.c/uns-outer-4.c b/libgomp/testsuite/libgomp.c/uns-outer-4.c
new file mode 100644
index 00000000000..cd646a54133
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/uns-outer-4.c
@@ -0,0 +1,36 @@
+/* { dg-do run } */
+/* { dg-additional-options "-ftree-parallelize-loops=2" } */
+
+void abort (void);
+
+unsigned int g_sum = 1;
+
+unsigned int x[500][500];
+
+void __attribute__((noinline,noclone))
+parloop (int N)
+{
+  int i, j;
+  unsigned int sum;
+
+  /* Double reduction is detected, outer loop is parallelized.  */
+  sum = 0;
+  for (i = 0; i < N; i++)
+    for (j = 0; j < N; j++)
+      sum += x[i][j];
+
+  g_sum = sum;
+}
+
+int
+main (void)
+{
+  x[234][432] = 2;
+
+  parloop (500);
+
+  if (g_sum != 2)
+    abort ();
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c++/c++.exp b/libgomp/testsuite/libgomp.oacc-c++/c++.exp
index 3b97024492c..88b0269e3f3 100644
--- a/libgomp/testsuite/libgomp.oacc-c++/c++.exp
+++ b/libgomp/testsuite/libgomp.oacc-c++/c++.exp
@@ -81,9 +81,6 @@ if { $lang_test_file_found } {
 	    host {
 		set acc_mem_shared 1
 	    }
-	    host_nonshm {
-		set acc_mem_shared 0
-	    }
 	    nvidia {
 		if { ![check_effective_target_openacc_nvidia_accel_present] } {
 		    # Don't bother; execution testing is going to FAIL.
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_on_device-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_on_device-1.c
index 81ea47689a9..8112745bcb8 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_on_device-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_on_device-1.c
@@ -15,8 +15,6 @@ main (int argc, char *argv[])
       abort ();
     if (!acc_on_device (acc_device_host))
       abort ();
-    if (acc_on_device (acc_device_host_nonshm))
-      abort ();
     if (acc_on_device (acc_device_not_host))
       abort ();
     if (acc_on_device (acc_device_nvidia))
@@ -32,8 +30,6 @@ main (int argc, char *argv[])
       abort ();
     if (!acc_on_device (acc_device_host))
       abort ();
-    if (acc_on_device (acc_device_host_nonshm))
-      abort ();
     if (acc_on_device (acc_device_not_host))
       abort ();
     if (acc_on_device (acc_device_nvidia))
@@ -51,13 +47,6 @@ main (int argc, char *argv[])
       abort ();
     if (acc_on_device (acc_device_host))
       abort ();
-#if ACC_DEVICE_TYPE_host_nonshm
-    if (!acc_on_device (acc_device_host_nonshm))
-      abort ();
-#else
-    if (acc_on_device (acc_device_host_nonshm))
-      abort ();
-#endif
     if (!acc_on_device (acc_device_not_host))
       abort ();
 #if ACC_DEVICE_TYPE_nvidia
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/if-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/if-1.c
index 184b3554bf5..6aa3bb7199e 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/if-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/if-1.c
@@ -1,5 +1,4 @@
 /* { dg-do run } */
-/* { dg-additional-options "-fno-builtin-acc_on_device" } */
 
 #include <openacc.h>
 #include <stdlib.h>
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-type-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-type-1.c
new file mode 100644
index 00000000000..5adfcecd641
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-type-1.c
@@ -0,0 +1,24 @@
+#define vector __attribute__ ((vector_size (4 * sizeof(int))))
+
+int main(void)
+{
+  vector int vi = { 12, -34, -56, 78 };
+
+#pragma acc parallel copy(vi)
+  {
+    if (vi[0] != 12
+	|| vi[1] != -34
+	|| vi[2] != -56
+	|| vi[3] != 78)
+      __builtin_abort();
+    vector int vi_ = { -21, -43, 65, 87 };
+    vi = vi_;
+  }
+  if (vi[0] != -21
+      || vi[1] != -43
+      || vi[2] != 65
+      || vi[3] != 87)
+    __builtin_abort();
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c/c.exp b/libgomp/testsuite/libgomp.oacc-c/c.exp
index 326b988d99a..5020e6a0691 100644
--- a/libgomp/testsuite/libgomp.oacc-c/c.exp
+++ b/libgomp/testsuite/libgomp.oacc-c/c.exp
@@ -44,9 +44,6 @@ foreach offload_target_openacc $offload_targets_s_openacc {
 	host {
 	    set acc_mem_shared 1
 	}
-	host_nonshm {
-	    set acc_mem_shared 0
-	}
 	nvidia {
 	    if { ![check_effective_target_openacc_nvidia_accel_present] } {
 		# Don't bother; execution testing is going to FAIL.
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-1.f90
index 448881837ac..1a10f32ab3a 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-1.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-1.f90
@@ -11,7 +11,6 @@ implicit none
 
 if (.not. acc_on_device (acc_device_none)) call abort
 if (.not. acc_on_device (acc_device_host)) call abort
-if (acc_on_device (acc_device_host_nonshm)) call abort
 if (acc_on_device (acc_device_not_host)) call abort
 if (acc_on_device (acc_device_nvidia)) call abort
 
@@ -21,7 +20,6 @@ if (acc_on_device (acc_device_nvidia)) call abort
 !$acc parallel if(.false.)
 if (.not. acc_on_device (acc_device_none)) call abort
 if (.not. acc_on_device (acc_device_host)) call abort
-if (acc_on_device (acc_device_host_nonshm)) call abort
 if (acc_on_device (acc_device_not_host)) call abort
 if (acc_on_device (acc_device_nvidia)) call abort
 !$acc end parallel
@@ -34,11 +32,6 @@ if (acc_on_device (acc_device_nvidia)) call abort
 !$acc parallel
 if (acc_on_device (acc_device_none)) call abort
 if (acc_on_device (acc_device_host)) call abort
-#if ACC_DEVICE_TYPE_host_nonshm
-if (.not. acc_on_device (acc_device_host_nonshm)) call abort
-#else
-if (acc_on_device (acc_device_host_nonshm)) call abort
-#endif
 if (.not. acc_on_device (acc_device_not_host)) call abort
 #if ACC_DEVICE_TYPE_nvidia
 if (.not. acc_on_device (acc_device_nvidia)) call abort
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-2.f b/libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-2.f
index 0047a194f66..a19045b0563 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-2.f
+++ b/libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-2.f
@@ -11,7 +11,6 @@
 
       IF (.NOT. ACC_ON_DEVICE (ACC_DEVICE_NONE)) CALL ABORT
       IF (.NOT. ACC_ON_DEVICE (ACC_DEVICE_HOST)) CALL ABORT
-      IF (ACC_ON_DEVICE (ACC_DEVICE_HOST_NONSHM)) CALL ABORT
       IF (ACC_ON_DEVICE (ACC_DEVICE_NOT_HOST)) CALL ABORT
       IF (ACC_ON_DEVICE (ACC_DEVICE_NVIDIA)) CALL ABORT
 
@@ -21,7 +20,6 @@
 !$ACC PARALLEL IF(.FALSE.)
       IF (.NOT. ACC_ON_DEVICE (ACC_DEVICE_NONE)) CALL ABORT
       IF (.NOT. ACC_ON_DEVICE (ACC_DEVICE_HOST)) CALL ABORT
-      IF (ACC_ON_DEVICE (ACC_DEVICE_HOST_NONSHM)) CALL ABORT
       IF (ACC_ON_DEVICE (ACC_DEVICE_NOT_HOST)) CALL ABORT
       IF (ACC_ON_DEVICE (ACC_DEVICE_NVIDIA)) CALL ABORT
 !$ACC END PARALLEL
@@ -34,11 +32,6 @@
 !$ACC PARALLEL
       IF (ACC_ON_DEVICE (ACC_DEVICE_NONE)) CALL ABORT
       IF (ACC_ON_DEVICE (ACC_DEVICE_HOST)) CALL ABORT
-#if ACC_DEVICE_TYPE_host_nonshm
-      IF (.NOT. ACC_ON_DEVICE (ACC_DEVICE_HOST_NONSHM)) CALL ABORT
-#else
-      IF (ACC_ON_DEVICE (ACC_DEVICE_HOST_NONSHM)) CALL ABORT
-#endif
       IF (.NOT. ACC_ON_DEVICE (ACC_DEVICE_NOT_HOST)) CALL ABORT
 #if ACC_DEVICE_TYPE_nvidia
       IF (.NOT. ACC_ON_DEVICE (ACC_DEVICE_NVIDIA)) CALL ABORT
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-3.f b/libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-3.f
index 49d7a720fe3..c3917760682 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-3.f
+++ b/libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-3.f
@@ -11,7 +11,6 @@
 
       IF (.NOT. ACC_ON_DEVICE (ACC_DEVICE_NONE)) CALL ABORT
       IF (.NOT. ACC_ON_DEVICE (ACC_DEVICE_HOST)) CALL ABORT
-      IF (ACC_ON_DEVICE (ACC_DEVICE_HOST_NONSHM)) CALL ABORT
       IF (ACC_ON_DEVICE (ACC_DEVICE_NOT_HOST)) CALL ABORT
       IF (ACC_ON_DEVICE (ACC_DEVICE_NVIDIA)) CALL ABORT
 
@@ -21,7 +20,6 @@
 !$ACC PARALLEL IF(.FALSE.)
       IF (.NOT. ACC_ON_DEVICE (ACC_DEVICE_NONE)) CALL ABORT
       IF (.NOT. ACC_ON_DEVICE (ACC_DEVICE_HOST)) CALL ABORT
-      IF (ACC_ON_DEVICE (ACC_DEVICE_HOST_NONSHM)) CALL ABORT
       IF (ACC_ON_DEVICE (ACC_DEVICE_NOT_HOST)) CALL ABORT
       IF (ACC_ON_DEVICE (ACC_DEVICE_NVIDIA)) CALL ABORT
 !$ACC END PARALLEL
@@ -34,11 +32,6 @@
 !$ACC PARALLEL
       IF (ACC_ON_DEVICE (ACC_DEVICE_NONE)) CALL ABORT
       IF (ACC_ON_DEVICE (ACC_DEVICE_HOST)) CALL ABORT
-#if ACC_DEVICE_TYPE_host_nonshm
-      IF (.NOT. ACC_ON_DEVICE (ACC_DEVICE_HOST_NONSHM)) CALL ABORT
-#else
-      IF (ACC_ON_DEVICE (ACC_DEVICE_HOST_NONSHM)) CALL ABORT
-#endif
       IF (.NOT. ACC_ON_DEVICE (ACC_DEVICE_NOT_HOST)) CALL ABORT
 #if ACC_DEVICE_TYPE_nvidia
       IF (.NOT. ACC_ON_DEVICE (ACC_DEVICE_NVIDIA)) CALL ABORT
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/fortran.exp b/libgomp/testsuite/libgomp.oacc-fortran/fortran.exp
index a8aaff0e2b1..2d6b647af22 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/fortran.exp
+++ b/libgomp/testsuite/libgomp.oacc-fortran/fortran.exp
@@ -73,9 +73,6 @@ if { $lang_test_file_found } {
 	    host {
 		set acc_mem_shared 1
 	    }
-	    host_nonshm {
-		set acc_mem_shared 0
-	    }
 	    nvidia {
 		if { ![check_effective_target_openacc_nvidia_accel_present] } {
 		    # Don't bother; execution testing is going to FAIL.