This is the mail archive of the gcc-patches@gcc.gnu.org mailing list for the GCC project.
| Index Nav: | [Date Index] [Subject Index] [Author Index] [Thread Index] | |
|---|---|---|
| Message Nav: | [Date Prev] [Date Next] | [Thread Prev] [Thread Next] |
| Other format: | [Raw text] | |
Hi!
Resending with the auto-generated changes snipped; ezmlm thought this was
too big, and refused to deliver it to the mailing list.
On Wed, 17 Dec 2014 23:24:17 +0100, I wrote:
> Committed to gomp-4_0-branch in r218839:
>
> commit 1c4f05a68c6d0d5b6137bb6d85a293d16727b389
> Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
> Date: Wed Dec 17 22:23:02 2014 +0000
>
> libgomp updates.
>
> libgomp/
> * configure.ac: Rename from GNU OpenMP Runtime Library to GNU
> Offloading and Multi Processing Runtime Library. Change all
> users.
>
> libgomp/
> * plugin/Makefile.am: Rework into...
> * plugin/Makefrag.am: ... this new file, and
> * plugin/configure.ac: rework into...
> * plugin/configfrag.ac: ... this new file. Change all users.
> * configure.ac: Move plugin/offloading handling...
> * plugin/configfrag.ac: ... here.
>
> libgomp/
> * env.c (initialize_env): Don't look for GCC_ACC_NOTIFY but for
> GOMP_DEBUG. Document it in the libgomp manual.
> * libgomp-plugin.h: Rename GOMP_PLUGIN_notify to
> GOMP_PLUGIN_debug. Change all users.
> * libgomp.h: Rename goacc_notify_var to gomp_debug_var,
> gomp_vnotify to gomp_vdebug, and gomp_notify to gomp_debug.
> Change all users. Add kind argument to gomp_vdebug and
> gomp_debug. Change all users.
>
> libgomp/
> * libgomp_g.h: Remove names of formal parameters.
>
> libgomp/
> * libgomp_target.h: Rename ACC_dispatch_t to acc_dispatch_t.
> Change all users.
> * oacc-init.c: Rename _acc_init to acc_init_1, and _acc_shutdown
> to acc_shutdown_1. Change all users.
> * oacc-int.h: Rename ACC_register to goacc_register,
> ACC_runtime_initialize to goacc_runtime_initialize,
> ACC_save_and_set_bind to goacc_save_and_set_bind, ACC_restore_bind
> to goacc_restore_bind, and ACC_lazy_initialize to
> goacc_lazy_initialize. Change all users.
> * plugin/plugin-nvptx.c: Rename cuErrorList to cuda_errlist,
> cuErrorMsg to cuda_error, cuSymNames to cuda_symnames, PTX_inited
> to ptx_inited, PTX_stream to ptx_stream, PTX_device to ptx_device,
> PTX_event to ptx_event, PTX_event_type to ptx_event_type, PTX_init
> to nvptx_init, , PTX_fini to nvptx_fini, PTX_open_device to
> nvptx_open_device, PTX_close_device to nvptx_close_device,
> PTX_get_num_devices to nvptx_get_num_devices, PTX_exec to
> nvptx_exec, PTX_alloc to nvptx_alloc, PTX_free to nvptx_free,
> PTX_host2dev to nvptx_host2dev, PTX_dev2host to nvptx_dev2host,
> PTX_set_async to nvptx_set_async, PTX_async_test to
> nvptx_async_test, PTX_async_test_all to nvptx_async_test_all,
> PTX_wait to nvptx_wait, PTX_wait_async to nvptx_wait_async,
> PTX_wait_all to nvptx_wait_all, PTX_wait_all_async to
> nvptx_wait_all_async, PTX_get_current_cuda_device to
> nvptx_get_current_cuda_device, PTX_get_current_cuda_context to
> nvptx_get_current_cuda_context, PTX_get_cuda_stream to
> nvptx_get_cuda_stream, PTX_set_cuda_stream to
> nvptx_set_cuda_stream. Change all users.
>
> libgomp/
> * oacc-parallel.c (GOACC_kernels): Pass 0 instead of num_waits in
> GOACC_parallel invocation.
>
> libgomp/
> * plugin/plugin-host.c: Remove all DEBUG code.
> * plugin/plugin-nvptx.c: Likewise.
>
> libgomp/
> * plugin/plugin-host.c (GOMP_OFFLOAD_get_caps): Don't claim
> TARGET_CAP_OPENMP_400.
>
> libgomp/
> * testsuite/libgomp.oacc-c-c++-common/lib-11.c: Restrict to target
> openacc_nvidia_accel_selected.
>
> libgomp/
> * testsuite/libgomp.oacc-c-c++-common/lib-38.c: Remove wrong
> check.
>
> libgomp/
> * testsuite/libgomp.oacc-c-c++-common/lib-9.c: Fix wrong check.
>
> include/
> * gomp-constants.h: Don't define GOMP_MAP_FORCE_PRIVATE and
> GOMP_MAP_FORCE_FIRSTPRIVATE. Change all users.
>
> git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@218839 138bc75d-0d04-0410-961f-82ee72b054a4
> ---
> contrib/gcc_update | 6 +-
> gcc/doc/install.texi | 3 +-
> gcc/doc/sourcebuild.texi | 2 +-
> gcc/fortran/gfortran.texi | 10 +-
> gcc/fortran/intrinsic.texi | 9 +-
> include/ChangeLog.gomp | 5 +
> include/gomp-constants.h | 5 +-
> libgomp/ChangeLog.gomp | 74 +
> libgomp/Makefile.am | 14 +-
> libgomp/Makefile.in | 125 +-
> libgomp/alloc.c | 3 +-
> libgomp/barrier.c | 3 +-
> libgomp/config.h.in | 9 +-
> libgomp/config/bsd/proc.c | 3 +-
> libgomp/config/linux/affinity.c | 3 +-
> libgomp/config/linux/alpha/futex.h | 3 +-
> libgomp/config/linux/bar.c | 3 +-
> libgomp/config/linux/bar.h | 3 +-
> libgomp/config/linux/futex.h | 3 +-
> libgomp/config/linux/ia64/futex.h | 3 +-
> libgomp/config/linux/lock.c | 3 +-
> libgomp/config/linux/mips/futex.h | 3 +-
> libgomp/config/linux/mutex.c | 3 +-
> libgomp/config/linux/mutex.h | 3 +-
> libgomp/config/linux/powerpc/futex.h | 3 +-
> libgomp/config/linux/proc.c | 3 +-
> libgomp/config/linux/proc.h | 3 +-
> libgomp/config/linux/ptrlock.c | 3 +-
> libgomp/config/linux/ptrlock.h | 3 +-
> libgomp/config/linux/s390/futex.h | 3 +-
> libgomp/config/linux/sem.c | 3 +-
> libgomp/config/linux/sem.h | 3 +-
> libgomp/config/linux/sparc/futex.h | 3 +-
> libgomp/config/linux/tile/futex.h | 3 +-
> libgomp/config/linux/wait.h | 3 +-
> libgomp/config/linux/x86/futex.h | 3 +-
> libgomp/config/mingw32/proc.c | 3 +-
> libgomp/config/mingw32/time.c | 3 +-
> libgomp/config/posix/affinity.c | 3 +-
> libgomp/config/posix/bar.c | 3 +-
> libgomp/config/posix/bar.h | 3 +-
> libgomp/config/posix/lock.c | 3 +-
> libgomp/config/posix/mutex.h | 3 +-
> libgomp/config/posix/proc.c | 3 +-
> libgomp/config/posix/ptrlock.h | 3 +-
> libgomp/config/posix/sem.c | 3 +-
> libgomp/config/posix/sem.h | 3 +-
> libgomp/config/posix/time.c | 3 +-
> libgomp/configure | 417 +-
> libgomp/configure.ac | 53 +-
> libgomp/critical.c | 3 +-
> libgomp/env.c | 10 +-
> libgomp/error.c | 43 +-
> libgomp/fortran.c | 3 +-
> libgomp/iter.c | 3 +-
> libgomp/iter_ull.c | 3 +-
> libgomp/libgomp-plugin.c | 24 +-
> libgomp/libgomp-plugin.h | 21 +-
> libgomp/libgomp.h | 31 +-
> libgomp/libgomp.map | 2 +-
> libgomp/libgomp.texi | 43 +-
> libgomp/libgomp_f.h.in | 3 +-
> libgomp/libgomp_g.h | 33 +-
> libgomp/libgomp_target.h | 9 +-
> libgomp/loop.c | 3 +-
> libgomp/loop_ull.c | 3 +-
> libgomp/oacc-async.c | 3 +-
> libgomp/oacc-cuda.c | 5 +-
> libgomp/oacc-host.c | 5 +-
> libgomp/oacc-init.c | 29 +-
> libgomp/oacc-int.h | 13 +-
> libgomp/oacc-mem.c | 23 +-
> libgomp/oacc-parallel.c | 79 +-
> libgomp/oacc-plugin.c | 3 +-
> libgomp/oacc-plugin.h | 3 +-
> libgomp/omp.h.in | 3 +-
> libgomp/omp_lib.f90.in | 3 +-
> libgomp/omp_lib.h.in | 3 +-
> libgomp/openacc.f90 | 3 +-
> libgomp/openacc.h | 9 +-
> libgomp/openacc_lib.h | 3 +-
> libgomp/ordered.c | 3 +-
> libgomp/parallel.c | 3 +-
> libgomp/plugin/Makefile.in | 630 -
> libgomp/plugin/{Makefile.am => Makefrag.am} | 26 +-
> libgomp/plugin/aclocal.m4 | 978 --
> libgomp/plugin/config.h.in | 65 -
> libgomp/plugin/configfrag.ac | 148 +
> libgomp/plugin/configure | 16835 -------------------
> libgomp/plugin/configure.ac | 179 -
> libgomp/plugin/plugin-host.c | 109 +-
> libgomp/plugin/plugin-nvptx.c | 756 +-
> libgomp/sections.c | 3 +-
> libgomp/single.c | 3 +-
> libgomp/splay-tree.c | 3 +-
> libgomp/splay-tree.h | 3 +-
> libgomp/target.c | 65 +-
> libgomp/task.c | 3 +-
> libgomp/team.c | 3 +-
> libgomp/testsuite/Makefile.in | 17 +-
> libgomp/testsuite/lib/libgomp.exp | 4 +-
> .../libgomp-test-support.exp.in | 0
> libgomp/testsuite/libgomp.oacc-c++/c++.exp | 4 +-
> libgomp/testsuite/libgomp.oacc-c-c++-common/if-1.c | 3 -
> .../testsuite/libgomp.oacc-c-c++-common/lib-11.c | 3 +-
> .../testsuite/libgomp.oacc-c-c++-common/lib-38.c | 3 -
> .../testsuite/libgomp.oacc-c-c++-common/lib-9.c | 2 +-
> libgomp/work.c | 3 +-
> liboffloadmic/plugin/Makefile.am | 3 +-
> liboffloadmic/plugin/Makefile.in | 3 +-
> liboffloadmic/plugin/configure.ac | 3 +-
> liboffloadmic/plugin/libgomp-plugin-intelmic.cpp | 3 +-
> liboffloadmic/plugin/offload_target_main.cpp | 3 +-
> libstdc++-v3/doc/xml/manual/parallel_mode.xml | 4 +-
> 114 files changed, 1316 insertions(+), 19824 deletions(-)
>
> diff --git contrib/gcc_update contrib/gcc_update
> index a50dc8c..5ba3a05 100755
> --- contrib/gcc_update
> +++ contrib/gcc_update
> @@ -139,14 +139,12 @@ libcpp/aclocal.m4: libcpp/configure.ac
> libcpp/Makefile.in: libcpp/configure.ac libcpp/aclocal.m4
> libcpp/configure: libcpp/configure.ac libcpp/aclocal.m4
> libgomp/aclocal.m4: libgomp/configure.ac libgomp/acinclude.m4
> +libgomp/Makefile.am: libgomp/plugin/Makefrag.am
> libgomp/Makefile.in: libgomp/Makefile.am libgomp/aclocal.m4
> libgomp/testsuite/Makefile.in: libgomp/testsuite/Makefile.am libgomp/aclocal.m4
> +libgomp/configure.ac: libgomp/plugin/configfrag.ac
> libgomp/configure: libgomp/configure.ac libgomp/aclocal.m4
> libgomp/config.h.in: libgomp/configure.ac libgomp/aclocal.m4
> -libgomp/plugin/aclocal.m4: libgomp/plugin/configure.ac
> -libgomp/plugin/Makefile.in: libgomp/plugin/Makefile.am libgomp/plugin/aclocal.m4
> -libgomp/plugin/configure: libgomp/plugin/configure.ac libgomp/plugin/aclocal.m4
> -libgomp/plugin/config.h.in: libgomp/plugin/configure.ac libgomp/plugin/aclocal.m4
> libitm/aclocal.m4: libitm/configure.ac libitm/acinclude.m4
> libitm/Makefile.in: libitm/Makefile.am libitm/aclocal.m4
> libitm/testsuite/Makefile.in: libitm/testsuite/Makefile.am libitm/aclocal.m4
> diff --git gcc/doc/install.texi gcc/doc/install.texi
> index 488c1f8..1a7cdd6 100644
> --- gcc/doc/install.texi
> +++ gcc/doc/install.texi
> @@ -1594,7 +1594,8 @@ Specify that the Fortran front end and @code{libgfortran} do not add
> support for @code{libquadmath} on systems supporting it.
>
> @item --disable-libgomp
> -Specify that the run-time libraries used by GOMP should not be built.
> +Specify that the GNU Offloading and Multi Processing Runtime Library
> +should not be built.
>
> @item --disable-libvtv
> Specify that the run-time libraries used by vtable verification
> diff --git gcc/doc/sourcebuild.texi gcc/doc/sourcebuild.texi
> index 661c5c9..724334c 100644
> --- gcc/doc/sourcebuild.texi
> +++ gcc/doc/sourcebuild.texi
> @@ -89,7 +89,7 @@ The Go runtime library. The bulk of this library is mirrored from the
> @uref{http://code.google.com/@/p/@/go/, master Go repository}.
>
> @item libgomp
> -The GNU OpenACC and OpenMP runtime library.
> +The GNU Offloading and Multi Processing Runtime Library.
>
> @item libiberty
> The @code{libiberty} library, used for portability and for some
> diff --git gcc/fortran/gfortran.texi gcc/fortran/gfortran.texi
> index 72540c3..63957af 100644
> --- gcc/fortran/gfortran.texi
> +++ gcc/fortran/gfortran.texi
> @@ -1912,8 +1912,9 @@ directives in fixed form; the @code{!$} conditional compilation
> sentinels in free form; and the @code{c$}, @code{*$} and @code{!$}
> sentinels in fixed form, @command{gfortran} needs to be invoked with
> the @option{-fopenacc}. This also arranges for automatic linking of
> -the GNU OpenACC runtime library @ref{Top,,libgomp,libgomp,GNU OpenACC
> -and OpenMP runtime library}.
> +the GNU Offloading and Multi Processing Runtime Library
> +@ref{Top,,libgomp,libgomp,GNU Offloading and Multi Processing Runtime
> +Library}.
>
> The OpenACC Fortran runtime library routines are provided both in a
> form of a Fortran 90 module named @code{openacc} and in a form of a
> @@ -1940,8 +1941,9 @@ directives in fixed form; the @code{!$} conditional compilation sentinels
> in free form; and the @code{c$}, @code{*$} and @code{!$} sentinels
> in fixed form, @command{gfortran} needs to be invoked with the
> @option{-fopenmp}. This also arranges for automatic linking of the
> -GNU OpenMP runtime library @ref{Top,,libgomp,libgomp,GNU OpenACC and OpenMP
> -runtime library}.
> +GNU Offloading and Multi Processing Runtime Library
> +@ref{Top,,libgomp,libgomp,GNU Offloading and Multi Processing Runtime
> +Library}.
>
> The OpenMP Fortran runtime library routines are provided both in a
> form of a Fortran 90 module named @code{omp_lib} and in a form of
> diff --git gcc/fortran/intrinsic.texi gcc/fortran/intrinsic.texi
> index fdaf044..bcce4ae 100644
> --- gcc/fortran/intrinsic.texi
> +++ gcc/fortran/intrinsic.texi
> @@ -14032,8 +14032,9 @@ The OpenACC Fortran runtime library routines are provided both in a
> form of a Fortran 90 module, named @code{OPENACC}, and in form of a
> Fortran @code{include} file named @file{openacc_lib.h}. The
> procedures provided by @code{OPENACC} can be found in the
> -@ref{Top,,Introduction,libgomp,GNU OpenACC and OpenMP runtime library}
> -manual, the named constants defined in the modules are listed below.
> +@ref{Top,,Introduction,libgomp,GNU Offloading and Multi Processing
> +Runtime Library} manual, the named constants defined in the modules
> +are listed below.
>
> For details refer to the actual
> @uref{http://www.openacc.org/,
> @@ -14058,8 +14059,8 @@ The OpenMP Fortran runtime library routines are provided both in
> a form of two Fortran 90 modules, named @code{OMP_LIB} and
> @code{OMP_LIB_KINDS}, and in a form of a Fortran @code{include} file named
> @file{omp_lib.h}. The procedures provided by @code{OMP_LIB} can be found
> -in the @ref{Top,,Introduction,libgomp,GNU OpenACC and OpenMP runtime
> -library} manual,
> +in the @ref{Top,,Introduction,libgomp,GNU Offloading and Multi
> +Processing Runtime Library} manual,
> the named constants defined in the modules are listed
> below.
>
> diff --git include/ChangeLog.gomp include/ChangeLog.gomp
> index 9172c26..6baa7a8 100644
> --- include/ChangeLog.gomp
> +++ include/ChangeLog.gomp
> @@ -1,3 +1,8 @@
> +2014-12-17 Thomas Schwinge <thomas@codesourcery.com>
> +
> + * gomp-constants.h: Don't define GOMP_MAP_FORCE_PRIVATE and
> + GOMP_MAP_FORCE_FIRSTPRIVATE. Change all users.
> +
> 2014-11-13 Thomas Schwinge <thomas@codesourcery.com>
>
> * gomp-constants.h: Define _GOMP_MAP_FLAG_SPECIAL and
> diff --git include/gomp-constants.h include/gomp-constants.h
> index 15b658f..e042f9e 100644
> --- include/gomp-constants.h
> +++ include/gomp-constants.h
> @@ -1,7 +1,8 @@
> /* Copyright (C) 2014 Free Software Foundation, Inc.
> Contributed by Mentor Embedded.
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> @@ -44,8 +45,6 @@
> #define GOMP_MAP_FORCE_PRESENT 0x0c
> #define GOMP_MAP_FORCE_DEALLOC 0x0d
> #define GOMP_MAP_FORCE_DEVICEPTR 0x0e
> -#define GOMP_MAP_FORCE_PRIVATE 0x18
> -#define GOMP_MAP_FORCE_FIRSTPRIVATE 0x19
>
> #define GOMP_MAP_COPYTO_P(X) \
> ((X) == GOMP_MAP_ALLOC_TO || (X) == GOMP_MAP_FORCE_TO)
> diff --git libgomp/ChangeLog.gomp libgomp/ChangeLog.gomp
> index b2c2526..ab668a5 100644
> --- libgomp/ChangeLog.gomp
> +++ libgomp/ChangeLog.gomp
> @@ -1,3 +1,77 @@
> +2014-12-17 Thomas Schwinge <thomas@codesourcery.com>
> + Julian Brown <julian@codesourcery.com>
> + David Malcolm <dmalcolm@redhat.com>
> +
> + * configure.ac: Rename from GNU OpenMP Runtime Library to GNU
> + Offloading and Multi Processing Runtime Library. Change all
> + users.
> +
> +2014-12-17 Thomas Schwinge <thomas@codesourcery.com>
> + Julian Brown <julian@codesourcery.com>
> +
> + * plugin/Makefile.am: Rework into...
> + * plugin/Makefrag.am: ... this new file, and
> + * plugin/configure.ac: rework into...
> + * plugin/configfrag.ac: ... this new file. Change all users.
> + * configure.ac: Move plugin/offloading handling...
> + * plugin/configfrag.ac: ... here.
> +
> + * env.c (initialize_env): Don't look for GCC_ACC_NOTIFY but for
> + GOMP_DEBUG. Document it in the libgomp manual.
> + * libgomp-plugin.h: Rename GOMP_PLUGIN_notify to
> + GOMP_PLUGIN_debug. Change all users.
> + * libgomp.h: Rename goacc_notify_var to gomp_debug_var,
> + gomp_vnotify to gomp_vdebug, and gomp_notify to gomp_debug.
> + Change all users. Add kind argument to gomp_vdebug and
> + gomp_debug. Change all users.
> +
> + * libgomp_g.h: Remove names of formal parameters.
> +
> + * libgomp_target.h: Rename ACC_dispatch_t to acc_dispatch_t.
> + Change all users.
> + * oacc-init.c: Rename _acc_init to acc_init_1, and _acc_shutdown
> + to acc_shutdown_1. Change all users.
> + * oacc-int.h: Rename ACC_register to goacc_register,
> + ACC_runtime_initialize to goacc_runtime_initialize,
> + ACC_save_and_set_bind to goacc_save_and_set_bind, ACC_restore_bind
> + to goacc_restore_bind, and ACC_lazy_initialize to
> + goacc_lazy_initialize. Change all users.
> + * plugin/plugin-nvptx.c: Rename cuErrorList to cuda_errlist,
> + cuErrorMsg to cuda_error, cuSymNames to cuda_symnames, PTX_inited
> + to ptx_inited, PTX_stream to ptx_stream, PTX_device to ptx_device,
> + PTX_event to ptx_event, PTX_event_type to ptx_event_type, PTX_init
> + to nvptx_init, , PTX_fini to nvptx_fini, PTX_open_device to
> + nvptx_open_device, PTX_close_device to nvptx_close_device,
> + PTX_get_num_devices to nvptx_get_num_devices, PTX_exec to
> + nvptx_exec, PTX_alloc to nvptx_alloc, PTX_free to nvptx_free,
> + PTX_host2dev to nvptx_host2dev, PTX_dev2host to nvptx_dev2host,
> + PTX_set_async to nvptx_set_async, PTX_async_test to
> + nvptx_async_test, PTX_async_test_all to nvptx_async_test_all,
> + PTX_wait to nvptx_wait, PTX_wait_async to nvptx_wait_async,
> + PTX_wait_all to nvptx_wait_all, PTX_wait_all_async to
> + nvptx_wait_all_async, PTX_get_current_cuda_device to
> + nvptx_get_current_cuda_device, PTX_get_current_cuda_context to
> + nvptx_get_current_cuda_context, PTX_get_cuda_stream to
> + nvptx_get_cuda_stream, PTX_set_cuda_stream to
> + nvptx_set_cuda_stream. Change all users.
> +
> + * oacc-parallel.c (GOACC_kernels): Pass 0 instead of num_waits in
> + GOACC_parallel invocation.
> +
> + * plugin/plugin-host.c: Remove all DEBUG code.
> + * plugin/plugin-nvptx.c: Likewise.
> +
> + * plugin/plugin-host.c (GOMP_OFFLOAD_get_caps): Don't claim
> + TARGET_CAP_OPENMP_400.
> +
> + * testsuite/libgomp.oacc-c-c++-common/lib-11.c: Restrict to target
> + openacc_nvidia_accel_selected.
> +
> + * testsuite/libgomp.oacc-c-c++-common/lib-38.c: Remove wrong
> + check.
> +
> + * testsuite/libgomp.oacc-c-c++-common/lib-9.c: Fix wrong check.
> +
> 2014-11-14 Tom de Vries <tom@codesourcery.com>
>
> * Makefile.am: Add missing dependency "openacc.mod: openacc.lo".
> diff --git libgomp/Makefile.am libgomp/Makefile.am
> index e5411ff..01bb1ec 100644
> --- libgomp/Makefile.am
> +++ libgomp/Makefile.am
> @@ -1,21 +1,21 @@
> ## Process this file with automake to produce Makefile.in
>
> ACLOCAL_AMFLAGS = -I .. -I ../config
> -SUBDIRS = testsuite plugin
> -DIST_SUBDIRS = plugin
> +SUBDIRS = testsuite
>
> ## May be used by toolexeclibdir.
> gcc_version := $(shell cat $(top_srcdir)/../gcc/BASE-VER)
>
> config_path = @config_path@
> -search_path = $(addprefix $(top_srcdir)/config/, $(config_path)) $(top_srcdir)
> +search_path = $(addprefix $(top_srcdir)/config/, $(config_path)) $(top_srcdir) \
> + $(top_srcdir)/../include
>
> fincludedir = $(libdir)/gcc/$(target_alias)/$(gcc_version)/finclude
> libsubincludedir = $(libdir)/gcc/$(target_alias)/$(gcc_version)/include
>
> vpath % $(strip $(search_path))
>
> -AM_CPPFLAGS = $(addprefix -I, $(search_path)) -I$(top_srcdir)/../include
> +AM_CPPFLAGS = $(addprefix -I, $(search_path))
> AM_CFLAGS = $(XCFLAGS)
> AM_LDFLAGS = $(XLDFLAGS) $(SECTION_LDFLAGS) $(OPT_LDFLAGS)
>
> @@ -62,8 +62,10 @@ libgomp_la_SOURCES = alloc.c barrier.c critical.c env.c error.c iter.c \
> iter_ull.c loop.c loop_ull.c ordered.c parallel.c sections.c single.c \
> task.c team.c work.c lock.c mutex.c proc.c sem.c bar.c ptrlock.c \
> time.c fortran.c affinity.c target.c oacc-parallel.c splay-tree.c \
> - oacc-host.c oacc-init.c oacc-mem.c oacc-async.c \
> - oacc-plugin.c oacc-cuda.c libgomp-plugin.c
> + oacc-host.c oacc-init.c oacc-mem.c oacc-async.c oacc-plugin.c \
> + oacc-cuda.c libgomp-plugin.c
> +
> +include $(top_srcdir)/plugin/Makefrag.am
>
> if USE_FORTRAN
> libgomp_la_SOURCES += openacc.f90
> diff --git libgomp/Makefile.in libgomp/Makefile.in
> index 6ec8b26..2447498 100644
> --- libgomp/Makefile.in
> +++ libgomp/Makefile.in
> [...]
> diff --git libgomp/alloc.c libgomp/alloc.c
> index 0ce171a..df23aec 100644
> --- libgomp/alloc.c
> +++ libgomp/alloc.c
> @@ -1,7 +1,8 @@
> /* Copyright (C) 2005-2014 Free Software Foundation, Inc.
> Contributed by Richard Henderson <rth@redhat.com>.
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> diff --git libgomp/barrier.c libgomp/barrier.c
> index e4864f3..4bfc006 100644
> --- libgomp/barrier.c
> +++ libgomp/barrier.c
> @@ -1,7 +1,8 @@
> /* Copyright (C) 2005-2014 Free Software Foundation, Inc.
> Contributed by Richard Henderson <rth@redhat.com>.
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> diff --git libgomp/config.h.in libgomp/config.h.in
> index a5e27ca..02547b1 100644
> --- libgomp/config.h.in
> +++ libgomp/config.h.in
> [...]
> diff --git libgomp/config/bsd/proc.c libgomp/config/bsd/proc.c
> index b94d585..f734957 100644
> --- libgomp/config/bsd/proc.c
> +++ libgomp/config/bsd/proc.c
> @@ -1,7 +1,8 @@
> /* Copyright (C) 2005-2014 Free Software Foundation, Inc.
> Contributed by Richard Henderson <rth@redhat.com>.
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> diff --git libgomp/config/linux/affinity.c libgomp/config/linux/affinity.c
> index 2f8a300..13800c4 100644
> --- libgomp/config/linux/affinity.c
> +++ libgomp/config/linux/affinity.c
> @@ -1,7 +1,8 @@
> /* Copyright (C) 2006-2014 Free Software Foundation, Inc.
> Contributed by Jakub Jelinek <jakub@redhat.com>.
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> diff --git libgomp/config/linux/alpha/futex.h libgomp/config/linux/alpha/futex.h
> index 21cf3cc..73173dd 100644
> --- libgomp/config/linux/alpha/futex.h
> +++ libgomp/config/linux/alpha/futex.h
> @@ -1,7 +1,8 @@
> /* Copyright (C) 2005-2014 Free Software Foundation, Inc.
> Contributed by Richard Henderson <rth@redhat.com>.
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> diff --git libgomp/config/linux/bar.c libgomp/config/linux/bar.c
> index ea1e08b..b544288 100644
> --- libgomp/config/linux/bar.c
> +++ libgomp/config/linux/bar.c
> @@ -1,7 +1,8 @@
> /* Copyright (C) 2005-2014 Free Software Foundation, Inc.
> Contributed by Richard Henderson <rth@redhat.com>.
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> diff --git libgomp/config/linux/bar.h libgomp/config/linux/bar.h
> index 7adb45f..bf4df60 100644
> --- libgomp/config/linux/bar.h
> +++ libgomp/config/linux/bar.h
> @@ -1,7 +1,8 @@
> /* Copyright (C) 2005-2014 Free Software Foundation, Inc.
> Contributed by Richard Henderson <rth@redhat.com>.
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> diff --git libgomp/config/linux/futex.h libgomp/config/linux/futex.h
> index 63334c7..c728457 100644
> --- libgomp/config/linux/futex.h
> +++ libgomp/config/linux/futex.h
> @@ -1,7 +1,8 @@
> /* Copyright (C) 2010-2014 Free Software Foundation, Inc.
> Contributed by ARM Ltd.
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> diff --git libgomp/config/linux/ia64/futex.h libgomp/config/linux/ia64/futex.h
> index c1b9d42..a28a536 100644
> --- libgomp/config/linux/ia64/futex.h
> +++ libgomp/config/linux/ia64/futex.h
> @@ -1,7 +1,8 @@
> /* Copyright (C) 2005-2014 Free Software Foundation, Inc.
> Contributed by Richard Henderson <rth@redhat.com>.
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> diff --git libgomp/config/linux/lock.c libgomp/config/linux/lock.c
> index 46d70e3..9812ba6 100644
> --- libgomp/config/linux/lock.c
> +++ libgomp/config/linux/lock.c
> @@ -1,7 +1,8 @@
> /* Copyright (C) 2005-2014 Free Software Foundation, Inc.
> Contributed by Richard Henderson <rth@redhat.com>.
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> diff --git libgomp/config/linux/mips/futex.h libgomp/config/linux/mips/futex.h
> index 7961d32..ce4dcb7 100644
> --- libgomp/config/linux/mips/futex.h
> +++ libgomp/config/linux/mips/futex.h
> @@ -1,7 +1,8 @@
> /* Copyright (C) 2005-2014 Free Software Foundation, Inc.
> Contributed by Ilie Garbacea <ilie@mips.com>, Chao-ying Fu <fu@mips.com>.
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> diff --git libgomp/config/linux/mutex.c libgomp/config/linux/mutex.c
> index cd93cd2..beaeefc 100644
> --- libgomp/config/linux/mutex.c
> +++ libgomp/config/linux/mutex.c
> @@ -1,7 +1,8 @@
> /* Copyright (C) 2005-2014 Free Software Foundation, Inc.
> Contributed by Richard Henderson <rth@redhat.com>.
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> diff --git libgomp/config/linux/mutex.h libgomp/config/linux/mutex.h
> index b36d33a..6c79e96 100644
> --- libgomp/config/linux/mutex.h
> +++ libgomp/config/linux/mutex.h
> @@ -1,7 +1,8 @@
> /* Copyright (C) 2005-2014 Free Software Foundation, Inc.
> Contributed by Richard Henderson <rth@redhat.com>.
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> diff --git libgomp/config/linux/powerpc/futex.h libgomp/config/linux/powerpc/futex.h
> index 22eefa0..584cdfb 100644
> --- libgomp/config/linux/powerpc/futex.h
> +++ libgomp/config/linux/powerpc/futex.h
> @@ -1,7 +1,8 @@
> /* Copyright (C) 2005-2014 Free Software Foundation, Inc.
> Contributed by Richard Henderson <rth@redhat.com>.
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> diff --git libgomp/config/linux/proc.c libgomp/config/linux/proc.c
> index fa89f1d..cac8d32 100644
> --- libgomp/config/linux/proc.c
> +++ libgomp/config/linux/proc.c
> @@ -1,7 +1,8 @@
> /* Copyright (C) 2005-2014 Free Software Foundation, Inc.
> Contributed by Jakub Jelinek <jakub@redhat.com>.
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> diff --git libgomp/config/linux/proc.h libgomp/config/linux/proc.h
> index 4b68474..2de7960 100644
> --- libgomp/config/linux/proc.h
> +++ libgomp/config/linux/proc.h
> @@ -1,7 +1,8 @@
> /* Copyright (C) 2011-2014 Free Software Foundation, Inc.
> Contributed by Uros Bizjak <ubizjak@gmail.com>
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> diff --git libgomp/config/linux/ptrlock.c libgomp/config/linux/ptrlock.c
> index 5f573ec..95d7f69 100644
> --- libgomp/config/linux/ptrlock.c
> +++ libgomp/config/linux/ptrlock.c
> @@ -1,7 +1,8 @@
> /* Copyright (C) 2008-2014 Free Software Foundation, Inc.
> Contributed by Jakub Jelinek <jakub@redhat.com>.
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> diff --git libgomp/config/linux/ptrlock.h libgomp/config/linux/ptrlock.h
> index ca470b1..82c59ff 100644
> --- libgomp/config/linux/ptrlock.h
> +++ libgomp/config/linux/ptrlock.h
> @@ -1,7 +1,8 @@
> /* Copyright (C) 2008-2014 Free Software Foundation, Inc.
> Contributed by Jakub Jelinek <jakub@redhat.com>.
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> diff --git libgomp/config/linux/s390/futex.h libgomp/config/linux/s390/futex.h
> index 3dc264b..4a505fd 100644
> --- libgomp/config/linux/s390/futex.h
> +++ libgomp/config/linux/s390/futex.h
> @@ -1,7 +1,8 @@
> /* Copyright (C) 2005-2014 Free Software Foundation, Inc.
> Contributed by Jakub Jelinek <jakub@redhat.com>.
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> diff --git libgomp/config/linux/sem.c libgomp/config/linux/sem.c
> index 328ba89..885d04d 100644
> --- libgomp/config/linux/sem.c
> +++ libgomp/config/linux/sem.c
> @@ -1,7 +1,8 @@
> /* Copyright (C) 2005-2014 Free Software Foundation, Inc.
> Contributed by Richard Henderson <rth@redhat.com>.
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> diff --git libgomp/config/linux/sem.h libgomp/config/linux/sem.h
> index 5aa7bc2..f9a24cd 100644
> --- libgomp/config/linux/sem.h
> +++ libgomp/config/linux/sem.h
> @@ -1,7 +1,8 @@
> /* Copyright (C) 2005-2014 Free Software Foundation, Inc.
> Contributed by Richard Henderson <rth@redhat.com>.
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> diff --git libgomp/config/linux/sparc/futex.h libgomp/config/linux/sparc/futex.h
> index e7f94c3..41999d0 100644
> --- libgomp/config/linux/sparc/futex.h
> +++ libgomp/config/linux/sparc/futex.h
> @@ -1,7 +1,8 @@
> /* Copyright (C) 2005-2014 Free Software Foundation, Inc.
> Contributed by Jakub Jelinek <jakub@redhat.com>.
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> diff --git libgomp/config/linux/tile/futex.h libgomp/config/linux/tile/futex.h
> index 8ce362b..be438d7 100644
> --- libgomp/config/linux/tile/futex.h
> +++ libgomp/config/linux/tile/futex.h
> @@ -1,7 +1,8 @@
> /* Copyright (C) 2011-2014 Free Software Foundation, Inc.
> Contributed by Walter Lee (walt@tilera.com)
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> diff --git libgomp/config/linux/wait.h libgomp/config/linux/wait.h
> index 08141f7..adfb52b 100644
> --- libgomp/config/linux/wait.h
> +++ libgomp/config/linux/wait.h
> @@ -1,7 +1,8 @@
> /* Copyright (C) 2008-2014 Free Software Foundation, Inc.
> Contributed by Jakub Jelinek <jakub@redhat.com>.
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> diff --git libgomp/config/linux/x86/futex.h libgomp/config/linux/x86/futex.h
> index b822c67..5ce2f75 100644
> --- libgomp/config/linux/x86/futex.h
> +++ libgomp/config/linux/x86/futex.h
> @@ -1,7 +1,8 @@
> /* Copyright (C) 2005-2014 Free Software Foundation, Inc.
> Contributed by Richard Henderson <rth@redhat.com>.
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> diff --git libgomp/config/mingw32/proc.c libgomp/config/mingw32/proc.c
> index 7fe9b5c..ac27095 100644
> --- libgomp/config/mingw32/proc.c
> +++ libgomp/config/mingw32/proc.c
> @@ -1,7 +1,8 @@
> /* Copyright (C) 2007-2014 Free Software Foundation, Inc.
> Contributed by Danny Smith <dannysmith@users.sourceforge.net>
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> diff --git libgomp/config/mingw32/time.c libgomp/config/mingw32/time.c
> index 54811f4..a224adc 100644
> --- libgomp/config/mingw32/time.c
> +++ libgomp/config/mingw32/time.c
> @@ -1,7 +1,8 @@
> /* Copyright (C) 2006-2014 Free Software Foundation, Inc.
> Contributed by Francois-Xavier Coudert <coudert@clipper.ens.fr>
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> diff --git libgomp/config/posix/affinity.c libgomp/config/posix/affinity.c
> index a4b3267..b31127a 100644
> --- libgomp/config/posix/affinity.c
> +++ libgomp/config/posix/affinity.c
> @@ -1,7 +1,8 @@
> /* Copyright (C) 2006-2014 Free Software Foundation, Inc.
> Contributed by Jakub Jelinek <jakub@redhat.com>.
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> diff --git libgomp/config/posix/bar.c libgomp/config/posix/bar.c
> index f6a1bb1..b927e67 100644
> --- libgomp/config/posix/bar.c
> +++ libgomp/config/posix/bar.c
> @@ -1,7 +1,8 @@
> /* Copyright (C) 2005-2014 Free Software Foundation, Inc.
> Contributed by Richard Henderson <rth@redhat.com>.
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> diff --git libgomp/config/posix/bar.h libgomp/config/posix/bar.h
> index e36b403..d75d751 100644
> --- libgomp/config/posix/bar.h
> +++ libgomp/config/posix/bar.h
> @@ -1,7 +1,8 @@
> /* Copyright (C) 2005-2014 Free Software Foundation, Inc.
> Contributed by Richard Henderson <rth@redhat.com>.
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> diff --git libgomp/config/posix/lock.c libgomp/config/posix/lock.c
> index c65e041..88b0c8a 100644
> --- libgomp/config/posix/lock.c
> +++ libgomp/config/posix/lock.c
> @@ -1,7 +1,8 @@
> /* Copyright (C) 2005-2014 Free Software Foundation, Inc.
> Contributed by Richard Henderson <rth@redhat.com>.
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> diff --git libgomp/config/posix/mutex.h libgomp/config/posix/mutex.h
> index c0d6fd9..ede13e4 100644
> --- libgomp/config/posix/mutex.h
> +++ libgomp/config/posix/mutex.h
> @@ -1,7 +1,8 @@
> /* Copyright (C) 2005-2014 Free Software Foundation, Inc.
> Contributed by Richard Henderson <rth@redhat.com>.
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> diff --git libgomp/config/posix/proc.c libgomp/config/posix/proc.c
> index ded9c50..064ae64 100644
> --- libgomp/config/posix/proc.c
> +++ libgomp/config/posix/proc.c
> @@ -1,7 +1,8 @@
> /* Copyright (C) 2005-2014 Free Software Foundation, Inc.
> Contributed by Richard Henderson <rth@redhat.com>.
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> diff --git libgomp/config/posix/ptrlock.h libgomp/config/posix/ptrlock.h
> index 74d9b03..4255cfb 100644
> --- libgomp/config/posix/ptrlock.h
> +++ libgomp/config/posix/ptrlock.h
> @@ -1,7 +1,8 @@
> /* Copyright (C) 2008-2014 Free Software Foundation, Inc.
> Contributed by Jakub Jelinek <jakub@redhat.com>.
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> diff --git libgomp/config/posix/sem.c libgomp/config/posix/sem.c
> index c430e6a..73dae12 100644
> --- libgomp/config/posix/sem.c
> +++ libgomp/config/posix/sem.c
> @@ -1,7 +1,8 @@
> /* Copyright (C) 2005-2014 Free Software Foundation, Inc.
> Contributed by Richard Henderson <rth@redhat.com>.
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> diff --git libgomp/config/posix/sem.h libgomp/config/posix/sem.h
> index fb4345a..e1bb120 100644
> --- libgomp/config/posix/sem.h
> +++ libgomp/config/posix/sem.h
> @@ -1,7 +1,8 @@
> /* Copyright (C) 2005-2014 Free Software Foundation, Inc.
> Contributed by Richard Henderson <rth@redhat.com>.
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> diff --git libgomp/config/posix/time.c libgomp/config/posix/time.c
> index 6000aa3..7bbcaf6 100644
> --- libgomp/config/posix/time.c
> +++ libgomp/config/posix/time.c
> @@ -1,7 +1,8 @@
> /* Copyright (C) 2005-2014 Free Software Foundation, Inc.
> Contributed by Richard Henderson <rth@redhat.com>.
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> diff --git libgomp/configure libgomp/configure
> index cb9746d..f72378e 100755
> --- libgomp/configure
> +++ libgomp/configure
> [...]
> diff --git libgomp/configure.ac libgomp/configure.ac
> index 034c893..178c34d 100644
> --- libgomp/configure.ac
> +++ libgomp/configure.ac
> @@ -2,9 +2,7 @@
> # aclocal -I ../config && autoconf && autoheader && automake
>
> AC_PREREQ(2.64)
> -#TODO: Update for OpenACC? But then also have to update copyright notices in
> -#all source files...
> -AC_INIT([GNU OpenMP Runtime Library], 1.0,,[libgomp])
> +AC_INIT([GNU Offloading and Multi Processing Runtime Library], 1.0,,[libgomp])
> AC_CONFIG_HEADER(config.h)
>
> # -------
> @@ -195,16 +193,7 @@ AC_LINK_IFELSE(
> [],
> [AC_MSG_ERROR([Pthreads are required to build libgomp])])])
>
> -plugin_support=yes
> -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.])
> -elif test "x${enable_offload_targets-no}" != xno; then
> - AC_MSG_ERROR([Can't support offloading without support for plugins])
> -fi
> -
> -AC_CONFIG_SUBDIRS([plugin])
> +m4_include([plugin/configfrag.ac])
>
> # Check for functions needed.
> AC_CHECK_FUNCS(getloadavg clock_gettime strtoull)
> @@ -289,40 +278,6 @@ else
> multilib_arg=
> fi
>
> -# Get accel target and path to install tree of accel compiler
> -offload_additional_options=
> -offload_additional_lib_paths=
> -offload_targets=host_nonshm
> -if test x"$enable_offload_targets" != x; then
> - for tgt in `echo $enable_offload_targets | sed -e 's#,# #g'`; do
> - tgt_dir=`echo $tgt | grep '=' | sed 's/.*=//'`
> - tgt=`echo $tgt | sed 's/=.*//'`
> - case $tgt in
> - *-intelmic-* | *-intelmicemul-*)
> - tgt_name="intelmic" ;;
> - *)
> - AC_MSG_ERROR([unknown offload target specified]) ;;
> - esac
> - if test x"$offload_targets" = x; then
> - offload_targets=$tgt_name
> - else
> - offload_targets=$offload_targets,$tgt_name
> - fi
> - if test x"$tgt_dir" != x; then
> - offload_additional_options="$offload_additional_options -B$tgt_dir/libexec/gcc/\$(target_alias)/\$(gcc_version) -B$tgt_dir/bin"
> - offload_additional_lib_paths="$offload_additional_lib_paths:$tgt_dir/lib64:$tgt_dir/lib"
> - else
> - offload_additional_options="$offload_additional_options -B\$(libexecdir)/gcc/\$(target_alias)/\$(gcc_version) -B\$(bindir)"
> - offload_additional_lib_paths="$offload_additional_lib_paths:$toolexeclibdir"
> - fi
> - done
> -fi
> -AC_DEFINE_UNQUOTED(OFFLOAD_TARGETS, "$offload_targets",
> - [Define to hold the list of target names suitable for offloading.])
> -AC_SUBST(offload_targets)
> -AC_SUBST(offload_additional_options)
> -AC_SUBST(offload_additional_lib_paths)
> -
> # Set up the set of libraries that we need to link against for libgomp.
> # Note that the GOMP_SELF_SPEC in gcc.c may force -pthread,
> # which will force linkage against -lpthread (or equivalent for the system).
> @@ -396,6 +351,6 @@ AC_SUBST(OMP_NEST_LOCK_25_KIND)
> CFLAGS="$save_CFLAGS"
>
> AC_CONFIG_FILES(omp.h omp_lib.h omp_lib.f90 libgomp_f.h)
> -AC_CONFIG_FILES(Makefile testsuite/Makefile)
> -AC_CONFIG_FILES(libgomp.spec)
> +AC_CONFIG_FILES(Makefile testsuite/Makefile libgomp.spec)
> +AC_CONFIG_FILES([testsuite/libgomp-test-support.exp])
> AC_OUTPUT
> diff --git libgomp/critical.c libgomp/critical.c
> index 7051441..419f4c6 100644
> --- libgomp/critical.c
> +++ libgomp/critical.c
> @@ -1,7 +1,8 @@
> /* Copyright (C) 2005-2014 Free Software Foundation, Inc.
> Contributed by Richard Henderson <rth@redhat.com>.
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> diff --git libgomp/env.c libgomp/env.c
> index 26d2149..de33443 100644
> --- libgomp/env.c
> +++ libgomp/env.c
> @@ -1,7 +1,8 @@
> /* Copyright (C) 2005-2014 Free Software Foundation, Inc.
> Contributed by Richard Henderson <rth@redhat.com>.
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> @@ -78,8 +79,7 @@ char *gomp_bind_var_list;
> unsigned long gomp_bind_var_list_len;
> void **gomp_places_list;
> unsigned long gomp_places_list_len;
> -
> -int goacc_notify_var;
> +int gomp_debug_var;
> int goacc_device_num;
> char* goacc_device_type;
>
> @@ -1197,7 +1197,7 @@ initialize_env (void)
> gomp_global_icv.thread_limit_var
> = thread_limit_var > INT_MAX ? UINT_MAX : thread_limit_var;
> }
> - parse_int ("GCC_ACC_NOTIFY", &goacc_notify_var, true);
> + parse_int ("GOMP_DEBUG", &gomp_debug_var, true);
> #ifndef HAVE_SYNC_BUILTINS
> gomp_mutex_init (&gomp_managed_threads_lock);
> #endif
> @@ -1296,7 +1296,7 @@ initialize_env (void)
> goacc_parse_device_type ();
>
> /* Initialize OpenACC-specific internal state. */
> - ACC_runtime_initialize ();
> + goacc_runtime_initialize ();
> }
>
>
> diff --git libgomp/error.c libgomp/error.c
> index 320b4d2..2c175d0 100644
> --- libgomp/error.c
> +++ libgomp/error.c
> @@ -1,7 +1,8 @@
> /* Copyright (C) 2005-2014 Free Software Foundation, Inc.
> Contributed by Richard Henderson <rth@redhat.com>.
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> @@ -35,6 +36,25 @@
> #include <stdlib.h>
>
>
> +#undef gomp_vdebug
> +void
> +gomp_vdebug (int kind __attribute__((unused)), const char *msg, va_list list)
> +{
> + if (gomp_debug_var)
> + vfprintf (stderr, msg, list);
> +}
> +
> +#undef gomp_debug
> +void
> +gomp_debug (int kind, const char *msg, ...)
> +{
> + va_list list;
> +
> + va_start (list, msg);
> + gomp_vdebug (kind, msg, list);
> + va_end (list);
> +}
> +
> void
> gomp_verror (const char *fmt, va_list list)
> {
> @@ -68,25 +88,4 @@ gomp_fatal (const char *fmt, ...)
> va_start (list, fmt);
> gomp_vfatal (fmt, list);
> va_end (list);
> -
> - /* Unreachable. */
> - abort ();
> }
> -
> -void
> -gomp_vnotify (const char *msg, va_list list)
> -{
> - if (goacc_notify_var)
> - vfprintf (stderr, msg, list);
> -}
> -
> -void
> -gomp_notify (const char *msg, ...)
> -{
> - va_list list;
> -
> - va_start (list, msg);
> - gomp_vnotify (msg, list);
> - va_end (list);
> -}
> -
> diff --git libgomp/fortran.c libgomp/fortran.c
> index 1f30c51..c519086 100644
> --- libgomp/fortran.c
> +++ libgomp/fortran.c
> @@ -1,7 +1,8 @@
> /* Copyright (C) 2005-2014 Free Software Foundation, Inc.
> Contributed by Jakub Jelinek <jakub@redhat.com>.
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> diff --git libgomp/iter.c libgomp/iter.c
> index b2efd07..9e0e36b 100644
> --- libgomp/iter.c
> +++ libgomp/iter.c
> @@ -1,7 +1,8 @@
> /* Copyright (C) 2005-2014 Free Software Foundation, Inc.
> Contributed by Richard Henderson <rth@redhat.com>.
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> diff --git libgomp/iter_ull.c libgomp/iter_ull.c
> index d003dbb..a96fc3e 100644
> --- libgomp/iter_ull.c
> +++ libgomp/iter_ull.c
> @@ -1,7 +1,8 @@
> /* Copyright (C) 2005-2014 Free Software Foundation, Inc.
> Contributed by Richard Henderson <rth@redhat.com>.
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> diff --git libgomp/libgomp-plugin.c libgomp/libgomp-plugin.c
> index b5dfb37..b30c260 100644
> --- libgomp/libgomp-plugin.c
> +++ libgomp/libgomp-plugin.c
> @@ -1,7 +1,9 @@
> /* Copyright (C) 2014 Free Software Foundation, Inc.
> +
> Contributed by Mentor Embedded.
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> @@ -49,6 +51,16 @@ GOMP_PLUGIN_realloc (void *ptr, size_t size)
> }
>
> void
> +GOMP_PLUGIN_debug (int kind, const char *msg, ...)
> +{
> + va_list ap;
> +
> + va_start (ap, msg);
> + gomp_debug (kind, msg, ap);
> + va_end (ap);
> +}
> +
> +void
> GOMP_PLUGIN_error (const char *msg, ...)
> {
> va_list ap;
> @@ -59,16 +71,6 @@ GOMP_PLUGIN_error (const char *msg, ...)
> }
>
> void
> -GOMP_PLUGIN_notify (const char *msg, ...)
> -{
> - va_list ap;
> -
> - va_start (ap, msg);
> - gomp_vnotify (msg, ap);
> - va_end (ap);
> -}
> -
> -void
> GOMP_PLUGIN_fatal (const char *msg, ...)
> {
> va_list ap;
> diff --git libgomp/libgomp-plugin.h libgomp/libgomp-plugin.h
> index 1496437..aae9b5f 100644
> --- libgomp/libgomp-plugin.h
> +++ libgomp/libgomp-plugin.h
> @@ -1,7 +1,9 @@
> /* Copyright (C) 2014 Free Software Foundation, Inc.
> +
> Contributed by Mentor Embedded.
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> @@ -29,25 +31,20 @@
>
> #include "mutex.h"
>
> -/* alloc.c */
> -
> extern void *GOMP_PLUGIN_malloc (size_t) __attribute__((malloc));
> extern void *GOMP_PLUGIN_malloc_cleared (size_t) __attribute__((malloc));
> extern void *GOMP_PLUGIN_realloc (void *, size_t);
>
> -/* error.c */
> -
> -extern void GOMP_PLUGIN_notify(const char *msg, ...);
> +extern void GOMP_PLUGIN_debug (int, const char *, ...)
> + __attribute__((format (printf, 2, 3)));
> extern void GOMP_PLUGIN_error (const char *, ...)
> __attribute__((format (printf, 1, 2)));
> extern void GOMP_PLUGIN_fatal (const char *, ...)
> __attribute__((noreturn, format (printf, 1, 2)));
>
> -/* mutex.c */
> -
> -extern void GOMP_PLUGIN_mutex_init (gomp_mutex_t *mutex);
> -extern void GOMP_PLUGIN_mutex_destroy (gomp_mutex_t *mutex);
> -extern void GOMP_PLUGIN_mutex_lock (gomp_mutex_t *mutex);
> -extern void GOMP_PLUGIN_mutex_unlock (gomp_mutex_t *mutex);
> +extern void GOMP_PLUGIN_mutex_init (gomp_mutex_t *);
> +extern void GOMP_PLUGIN_mutex_destroy (gomp_mutex_t *);
> +extern void GOMP_PLUGIN_mutex_lock (gomp_mutex_t *);
> +extern void GOMP_PLUGIN_mutex_unlock (gomp_mutex_t *);
>
> #endif
> diff --git libgomp/libgomp.h libgomp/libgomp.h
> index 9a58065..dbf1628 100644
> --- libgomp/libgomp.h
> +++ libgomp/libgomp.h
> @@ -1,7 +1,8 @@
> /* Copyright (C) 2005-2014 Free Software Foundation, Inc.
> Contributed by Richard Henderson <rth@redhat.com>.
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> @@ -23,9 +24,10 @@
> <http://www.gnu.org/licenses/>. */
>
> /* This file contains data types and function declarations that are not
> - part of the official OpenMP user interface. There are declarations
> - in here that are part of the GNU OpenMP ABI, in that the compiler is
> - required to know about them and use them.
> + part of the official OpenACC or OpenMP user interfaces. There are
> + declarations in here that are part of the GNU Offloading and Multi
> + Processing ABI, in that the compiler is required to know about them
> + and use them.
>
> The convention is that the all caps prefix "GOMP" is used group items
> that are part of the external ABI, and the lower case prefix "gomp"
> @@ -255,8 +257,7 @@ extern char *gomp_bind_var_list;
> extern unsigned long gomp_bind_var_list_len;
> extern void **gomp_places_list;
> extern unsigned long gomp_places_list_len;
> -
> -extern int goacc_notify_var;
> +extern int gomp_debug_var;
> extern int goacc_device_num;
> extern char* goacc_device_type;
>
> @@ -538,12 +539,24 @@ extern void *gomp_realloc (void *, size_t);
>
> /* error.c */
>
> -extern void gomp_vnotify (const char *, va_list);
> -extern void gomp_notify(const char *msg, ...);
> +extern void gomp_vdebug (int, const char *, va_list);
> +extern void gomp_debug (int, const char *, ...)
> + __attribute__((format (printf, 2, 3)));
> +#define gomp_vdebug(KIND, FMT, VALIST) \
> + do { \
> + if (__builtin_expect (gomp_debug_var, 0)) \
> + (gomp_vdebug) ((KIND), (FMT), (VALIST)); \
> + } while (0)
> +#define gomp_debug(KIND, ...) \
> + do { \
> + if (__builtin_expect (gomp_debug_var, 0)) \
> + (gomp_debug) ((KIND), __VA_ARGS__); \
> + } while (0)
> extern void gomp_verror (const char *, va_list);
> extern void gomp_error (const char *, ...)
> __attribute__((format (printf, 1, 2)));
> -extern void gomp_vfatal (const char *, va_list);
> +extern void gomp_vfatal (const char *, va_list)
> + __attribute__((noreturn));
> extern void gomp_fatal (const char *, ...)
> __attribute__((noreturn, format (printf, 1, 2)));
>
> diff --git libgomp/libgomp.map libgomp/libgomp.map
> index 68dbd6d..aa1fdb8 100644
> --- libgomp/libgomp.map
> +++ libgomp/libgomp.map
> @@ -332,7 +332,7 @@ GOMP_PLUGIN_1.0 {
> GOMP_PLUGIN_malloc_cleared;
> GOMP_PLUGIN_realloc;
> GOMP_PLUGIN_error;
> - GOMP_PLUGIN_notify;
> + GOMP_PLUGIN_debug;
> GOMP_PLUGIN_fatal;
> GOMP_PLUGIN_mutex_init;
> GOMP_PLUGIN_mutex_destroy;
> diff --git libgomp/libgomp.texi libgomp/libgomp.texi
> index 6c2673b..9618d4c 100644
> --- libgomp/libgomp.texi
> +++ libgomp/libgomp.texi
> @@ -31,13 +31,13 @@ texts being (a) (see below), and with the Back-Cover Texts being (b)
> @ifinfo
> @dircategory GNU Libraries
> @direntry
> -* libgomp: (libgomp). GNU OpenACC and OpenMP runtime library
> +* libgomp: (libgomp). GNU Offloading and Multi Processing Runtime Library
> @end direntry
>
> -This manual documents the GNU implementation of the OpenACC API for
> -offloading of code to accelerator devices in C/C++ and Fortran and
> -the GNU implementation of the OpenMP API for
> -multi-platform shared-memory parallel programming in C/C++ and Fortran.
> +This manual documents libgomp, the GNU Offloading and Multi Processing
> +Runtime library. This is the GNU implementation of the OpenMP and
> +OpenACC APIs for parallel and accelerator programming in C/C++ and
> +Fortran.
>
> Published by the Free Software Foundation
> 51 Franklin Street, Fifth Floor
> @@ -71,13 +71,19 @@ Boston, MA 02110-1301, USA@*
> @top Introduction
> @cindex Introduction
>
> -This manual documents the usage of libgomp, the GNU implementation of the
> -@uref{http://www.openacc.org/, OpenACC} Application Programming Interface (API)
> -for offloading of code to accelerator devices in C/C++ and Fortran, and
> -the GNU implementation of the
> -@uref{http://www.openmp.org, OpenMP} Application Programming Interface (API)
> -for multi-platform shared-memory parallel programming in C/C++ and Fortran.
> +This manual documents the usage of libgomp, the GNU Offloading and
> +Multi Processing Runtime Library. This includes the GNU
> +implementation of the @uref{http://www.openmp.org, OpenMP} Application
> +Programming Interface (API) for multi-platform shared-memory parallel
> +programming in C/C++ and Fortran, and the GNU implementation of the
> +@uref{http://www.openacc.org/, OpenACC} Application Programming
> +Interface (API) for offloading of code to accelerator devices in C/C++
> +and Fortran.
>
> +Originally, libgomp implemented the GNU OpenMP Runtime Library. Based
> +on this, support for OpenACC and offloading (both OpenACC and OpenMP
> +4's target construct) has been added later on, and the library's name
> +changed to GNU Offloading and Multi Processing Runtime Library.
>
>
> @comment
> @@ -1909,6 +1915,7 @@ beginning with @env{GOMP_} are GNU extensions.
> * 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
> +* GOMP_DEBUG:: Enable debugging output
> * GOMP_STACKSIZE:: Set default thread stack size
> * GOMP_SPINCOUNT:: Set the busy-wait spin count
> @end menu
> @@ -2228,6 +2235,20 @@ If both @env{GOMP_CPU_AFFINITY} and @env{OMP_PROC_BIND} are set,
>
>
>
> +@node GOMP_DEBUG
> +@section @env{GOMP_DEBUG} -- Enable debugging output
> +@cindex Environment Variable
> +@table @asis
> +@item @emph{Description}:
> +Enable debugging output. The variable should be set to @code{0}
> +(disabled, also the default if not set), or @code{1} (enabled).
> +
> +If enabled, some debugging output will be printed during execution.
> +This is currently not specified in more detail, and subject to change.
> +@end table
> +
> +
> +
> @node GOMP_STACKSIZE
> @section @env{GOMP_STACKSIZE} -- Set default thread stack size
> @cindex Environment Variable
> diff --git libgomp/libgomp_f.h.in libgomp/libgomp_f.h.in
> index 7de9528..bcf697b 100644
> --- libgomp/libgomp_f.h.in
> +++ libgomp/libgomp_f.h.in
> @@ -1,7 +1,8 @@
> /* Copyright (C) 2005-2014 Free Software Foundation, Inc.
> Contributed by Jakub Jelinek <jakub@redhat.com>.
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> diff --git libgomp/libgomp_g.h libgomp/libgomp_g.h
> index 932f0e2..4206d51 100644
> --- libgomp/libgomp_g.h
> +++ libgomp/libgomp_g.h
> @@ -1,7 +1,8 @@
> /* Copyright (C) 2005-2014 Free Software Foundation, Inc.
> Contributed by Richard Henderson <rth@redhat.com>.
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> @@ -219,29 +220,23 @@ extern void GOMP_teams (unsigned int, unsigned int);
> extern void GOACC_data_start (int, const void *,
> size_t, void **, size_t *, unsigned short *);
> extern void GOACC_data_end (void);
> -extern void GOACC_enter_exit_data (int device, const void *openmp_target,
> - size_t mapnum, void **hostaddrs,
> - size_t *sizes, unsigned short *kinds,
> - int async, int num_waits, ...);
> -extern void GOACC_kernels (int, void (*) (void *), const void *,
> - size_t, void **, size_t *, unsigned short *,
> - int, int, int, int, int, ...);
> -extern void GOACC_parallel (int, void (*) (void *), const void *,
> - size_t, void **, size_t *, unsigned short *,
> - int, int, int, int, int, ...);
> -extern void GOACC_update (int device, const void *openmp_target, size_t mapnum,
> - void **hostaddrs, size_t *sizes,
> - unsigned short *kinds, int async,
> - int num_waits, ...);
> +extern void GOACC_enter_exit_data (int, const void *, size_t, void **,
> + size_t *, unsigned short *, int, int, ...);
> +extern void GOACC_kernels (int, void (*) (void *), const void *, size_t,
> + void **, size_t *, unsigned short *, int, int, int,
> + int, int, ...);
> +extern void GOACC_parallel (int, void (*) (void *), const void *, size_t,
> + void **, size_t *, unsigned short *, int, int, int,
> + int, int, ...);
> +extern void GOACC_update (int, const void *, size_t, void **, size_t *,
> + unsigned short *, int, int, ...);
> extern void GOACC_wait (int, int, ...);
> extern int GOACC_get_num_threads (void);
> extern int GOACC_get_thread_num (void);
>
> /* oacc-mem.c */
>
> -extern void gomp_acc_insert_pointer (size_t mapnum, void **hostaddrs,
> - size_t *sizes, void *kinds);
> -extern void gomp_acc_remove_pointer (void *h, bool force_copyfrom, int async,
> - int mapnum);
> +extern void gomp_acc_insert_pointer (size_t, void **, size_t *, void *);
> +extern void gomp_acc_remove_pointer (void *, bool, int, int);
>
> #endif /* LIBGOMP_G_H */
> diff --git libgomp/libgomp_target.h libgomp/libgomp_target.h
> index b780259..6da9be8 100644
> --- libgomp/libgomp_target.h
> +++ libgomp/libgomp_target.h
> @@ -1,6 +1,7 @@
> /* Copyright (C) 2014 Free Software Foundation, Inc.
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> @@ -93,7 +94,7 @@ struct gomp_memory_mapping
> bool is_initialized;
> };
>
> -typedef struct ACC_dispatch_t
> +typedef struct acc_dispatch_t
> {
> /* This is a linked list of data mapped using the
> acc_map_data/acc_unmap_data or "acc enter data"/"acc exit data" pragmas
> @@ -136,7 +137,7 @@ typedef struct ACC_dispatch_t
> void *(*get_stream_func) (int);
> int (*set_stream_func) (int, void *);
> } cuda;
> -} ACC_dispatch_t;
> +} acc_dispatch_t;
>
> /* This structure describes accelerator device.
> It contains name of the corresponding libgomp plugin, function handlers for
> @@ -182,7 +183,7 @@ struct gomp_device_descr
> void (*run_func) (int, void *, void *);
>
> /* OpenACC-specific functions. */
> - ACC_dispatch_t openacc;
> + acc_dispatch_t openacc;
>
> /* Memory-mapping info for this device instance. */
> struct gomp_memory_mapping mem_map;
> diff --git libgomp/loop.c libgomp/loop.c
> index 65f3fe8..54c7e66 100644
> --- libgomp/loop.c
> +++ libgomp/loop.c
> @@ -1,7 +1,8 @@
> /* Copyright (C) 2005-2014 Free Software Foundation, Inc.
> Contributed by Richard Henderson <rth@redhat.com>.
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> diff --git libgomp/loop_ull.c libgomp/loop_ull.c
> index 0264eee..92ee5d9 100644
> --- libgomp/loop_ull.c
> +++ libgomp/loop_ull.c
> @@ -1,7 +1,8 @@
> /* Copyright (C) 2005-2014 Free Software Foundation, Inc.
> Contributed by Richard Henderson <rth@redhat.com>.
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> diff --git libgomp/oacc-async.c libgomp/oacc-async.c
> index 94c62d8..ec711f1 100644
> --- libgomp/oacc-async.c
> +++ libgomp/oacc-async.c
> @@ -4,7 +4,8 @@
>
> Contributed by Mentor Embedded.
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> diff --git libgomp/oacc-cuda.c libgomp/oacc-cuda.c
> index 9965d5c..8a81f03 100644
> --- libgomp/oacc-cuda.c
> +++ libgomp/oacc-cuda.c
> @@ -4,7 +4,8 @@
>
> Contributed by Mentor Embedded.
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> @@ -75,7 +76,7 @@ acc_set_cuda_stream (int async, void *stream)
> if (async < 0 || stream == NULL)
> return 0;
>
> - ACC_lazy_initialize ();
> + goacc_lazy_initialize ();
>
> if (base_dev && base_dev->openacc.cuda.set_stream_func)
> s = base_dev->openacc.cuda.set_stream_func (async, stream);
> diff --git libgomp/oacc-host.c libgomp/oacc-host.c
> index e116c08..f1ec426 100644
> --- libgomp/oacc-host.c
> +++ libgomp/oacc-host.c
> @@ -4,7 +4,8 @@
>
> Contributed by Mentor Embedded.
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> @@ -95,5 +96,5 @@ static __attribute__ ((constructor))
> void goacc_host_init (void)
> {
> gomp_mutex_init (&host_dispatch.mem_map.lock);
> - ACC_register (&host_dispatch);
> + goacc_register (&host_dispatch);
> }
> diff --git libgomp/oacc-init.c libgomp/oacc-init.c
> index 3d29eb3..7298d9a 100644
> --- libgomp/oacc-init.c
> +++ libgomp/oacc-init.c
> @@ -4,7 +4,8 @@
>
> Contributed by Mentor Embedded.
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> @@ -66,7 +67,7 @@ static gomp_mutex_t goacc_thread_lock;
> static struct gomp_device_descr const *dispatchers[_ACC_device_hwm] = { 0 };
>
> attribute_hidden void
> -ACC_register (struct gomp_device_descr const *disp)
> +goacc_register (struct gomp_device_descr const *disp)
> {
> /* Only register the 0th device here. */
> if (disp->target_id != 0)
> @@ -142,10 +143,10 @@ resolve_device (acc_device_t d)
>
> /* This is called when plugins have been initialized, and serves to call
> (indirectly) the target's device_init hook. Calling multiple times without
> - an intervening _acc_shutdown call is an error. */
> + an intervening acc_shutdown_1 call is an error. */
>
> static struct gomp_device_descr const *
> -_acc_init (acc_device_t d)
> +acc_init_1 (acc_device_t d)
> {
> struct gomp_device_descr const *acc_dev;
>
> @@ -288,7 +289,7 @@ acc_init (acc_device_t d)
>
> gomp_mutex_lock (&acc_device_lock);
>
> - base_dev = _acc_init (d);
> + base_dev = acc_init_1 (d);
>
> lazy_open (-1);
>
> @@ -297,8 +298,8 @@ acc_init (acc_device_t d)
>
> ialias (acc_init)
>
> -void
> -_acc_shutdown (acc_device_t d)
> +static void
> +acc_shutdown_1 (acc_device_t d)
> {
> struct goacc_thread *walk;
>
> @@ -353,7 +354,7 @@ acc_shutdown (acc_device_t d)
> {
> gomp_mutex_lock (&acc_device_lock);
>
> - _acc_shutdown (d);
> + acc_shutdown_1 (d);
>
> gomp_mutex_unlock (&acc_device_lock);
> }
> @@ -376,12 +377,12 @@ lazy_init (acc_device_t d)
> if (d == init_key)
> return base_dev;
>
> - _acc_shutdown (init_key);
> + acc_shutdown_1 (init_key);
> }
>
> assert (!base_dev);
>
> - return _acc_init (d);
> + return acc_init_1 (d);
> }
>
> /* Ensure that plugins are loaded, initialize and open the (default-numbered)
> @@ -554,7 +555,7 @@ acc_on_device (acc_device_t dev)
> ialias (acc_on_device)
>
> attribute_hidden void
> -ACC_runtime_initialize (void)
> +goacc_runtime_initialize (void)
> {
> gomp_mutex_init (&acc_device_lock);
>
> @@ -573,7 +574,7 @@ ACC_runtime_initialize (void)
> /* Compiler helper functions */
>
> attribute_hidden void
> -ACC_save_and_set_bind (acc_device_t d)
> +goacc_save_and_set_bind (acc_device_t d)
> {
> struct goacc_thread *thr = goacc_thread ();
>
> @@ -584,7 +585,7 @@ ACC_save_and_set_bind (acc_device_t d)
> }
>
> attribute_hidden void
> -ACC_restore_bind (void)
> +goacc_restore_bind (void)
> {
> struct goacc_thread *thr = goacc_thread ();
>
> @@ -598,7 +599,7 @@ ACC_restore_bind (void)
> pointers will be valid. */
>
> attribute_hidden void
> -ACC_lazy_initialize (void)
> +goacc_lazy_initialize (void)
> {
> struct goacc_thread *thr = goacc_thread ();
>
> diff --git libgomp/oacc-int.h libgomp/oacc-int.h
> index e2b6f7c..3c2c37f 100644
> --- libgomp/oacc-int.h
> +++ libgomp/oacc-int.h
> @@ -4,7 +4,8 @@
>
> Contributed by Mentor Embedded.
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> @@ -89,15 +90,15 @@ goacc_thread (void)
>
> struct gomp_device_descr;
>
> -void ACC_register (struct gomp_device_descr const *) __GOACC_NOTHROW;
> +void goacc_register (struct gomp_device_descr const *) __GOACC_NOTHROW;
>
> /* Current dispatcher. */
> extern struct gomp_device_descr const *base_dev;
>
> -void ACC_runtime_initialize (void);
> -void ACC_save_and_set_bind (acc_device_t);
> -void ACC_restore_bind (void);
> -void ACC_lazy_initialize (void);
> +void goacc_runtime_initialize (void);
> +void goacc_save_and_set_bind (acc_device_t);
> +void goacc_restore_bind (void);
> +void goacc_lazy_initialize (void);
>
> #ifdef HAVE_ATTRIBUTE_VISIBILITY
> # pragma GCC visibility pop
> diff --git libgomp/oacc-mem.c libgomp/oacc-mem.c
> index a8c5569..cfb63f5 100644
> --- libgomp/oacc-mem.c
> +++ libgomp/oacc-mem.c
> @@ -4,7 +4,8 @@
>
> Contributed by Mentor Embedded.
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> @@ -110,7 +111,7 @@ acc_malloc (size_t s)
> if (!s)
> return NULL;
>
> - ACC_lazy_initialize ();
> + goacc_lazy_initialize ();
>
> struct goacc_thread *thr = goacc_thread ();
>
> @@ -174,7 +175,7 @@ acc_deviceptr (void *h)
> void *d;
> void *offset;
>
> - ACC_lazy_initialize ();
> + goacc_lazy_initialize ();
>
> struct goacc_thread *thr = goacc_thread ();
>
> @@ -200,7 +201,7 @@ acc_hostptr (void *d)
> void *h;
> void *offset;
>
> - ACC_lazy_initialize ();
> + goacc_lazy_initialize ();
>
> struct goacc_thread *thr = goacc_thread ();
>
> @@ -226,7 +227,7 @@ acc_is_present (void *h, size_t s)
> if (!s || !h)
> return 0;
>
> - ACC_lazy_initialize ();
> + goacc_lazy_initialize ();
>
> struct goacc_thread *thr = goacc_thread ();
> struct gomp_device_descr *acc_dev = thr->dev;
> @@ -253,7 +254,7 @@ acc_map_data (void *h, void *d, size_t s)
> size_t sizes = s;
> unsigned short kinds = GOMP_MAP_ALLOC;
>
> - ACC_lazy_initialize ();
> + goacc_lazy_initialize ();
>
> struct goacc_thread *thr = goacc_thread ();
> struct gomp_device_descr *acc_dev = thr->dev;
> @@ -355,7 +356,7 @@ present_create_copy (unsigned f, void *h, size_t s)
> if (!h || !s)
> gomp_fatal ("[%p,+%d] is a bad range", (void *)h, (int)s);
>
> - ACC_lazy_initialize ();
> + goacc_lazy_initialize ();
>
> struct goacc_thread *thr = goacc_thread ();
> struct gomp_device_descr *acc_dev = thr->dev;
> @@ -517,10 +518,10 @@ gomp_acc_insert_pointer (size_t mapnum, void **hostaddrs, size_t *sizes,
> struct goacc_thread *thr = goacc_thread ();
> struct gomp_device_descr *acc_dev = thr->dev;
>
> - gomp_notify (" %s: prepare mappings\n", __FUNCTION__);
> + gomp_debug (0, " %s: prepare mappings\n", __FUNCTION__);
> tgt = gomp_map_vars ((struct gomp_device_descr *) acc_dev, mapnum, hostaddrs,
> NULL, sizes, kinds, true, false);
> - gomp_notify (" %s: mappings prepared\n", __FUNCTION__);
> + gomp_debug (0, " %s: mappings prepared\n", __FUNCTION__);
> tgt->prev = acc_dev->openacc.data_environ;
> acc_dev->openacc.data_environ = tgt;
> }
> @@ -539,7 +540,7 @@ gomp_acc_remove_pointer (void *h, bool force_copyfrom, int async, int mapnum)
> if (!n)
> gomp_fatal ("%p is not a mapped block", (void *)h);
>
> - gomp_notify (" %s: restore mappings\n", __FUNCTION__);
> + gomp_debug (0, " %s: restore mappings\n", __FUNCTION__);
>
> t = n->tgt;
>
> @@ -583,5 +584,5 @@ gomp_acc_remove_pointer (void *h, bool force_copyfrom, int async, int mapnum)
> acc_dev->openacc.register_async_cleanup_func (t);
> }
>
> - gomp_notify (" %s: mappings restored\n", __FUNCTION__);
> + gomp_debug (0, " %s: mappings restored\n", __FUNCTION__);
> }
> diff --git libgomp/oacc-parallel.c libgomp/oacc-parallel.c
> index ff51808..3726794 100644
> --- libgomp/oacc-parallel.c
> +++ libgomp/oacc-parallel.c
> @@ -1,8 +1,9 @@
> /* Copyright (C) 2013-2014 Free Software Foundation, Inc.
>
> - Contributed by Thomas Schwinge <thomas@codesourcery.com>.
> + Contributed by Mentor Embedded.
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> @@ -40,36 +41,31 @@
> static void
> dump_var (char *s, size_t idx, void *hostaddr, size_t size, unsigned char kind)
> {
> - gomp_notify(" %2zi: %3s 0x%.2x -", idx, s, kind & 0xff);
> + gomp_debug (0, " %2zi: %3s 0x%.2x -", idx, s, kind & 0xff);
>
> switch (kind & 0xff)
> {
> - case 0x00: gomp_notify(" ALLOC "); break;
> - case 0x01: gomp_notify(" ALLOC TO "); break;
> - case 0x02: gomp_notify(" ALLOC FROM "); break;
> - case 0x03: gomp_notify(" ALLOC TOFROM "); break;
> - case 0x04: gomp_notify(" POINTER "); break;
> - case 0x05: gomp_notify(" TO_PSET "); break;
> + case 0x00: gomp_debug (0, " ALLOC "); break;
> + case 0x01: gomp_debug (0, " ALLOC TO "); break;
> + case 0x02: gomp_debug (0, " ALLOC FROM "); break;
> + case 0x03: gomp_debug (0, " ALLOC TOFROM "); break;
> + case 0x04: gomp_debug (0, " POINTER "); break;
> + case 0x05: gomp_debug (0, " TO_PSET "); break;
>
> - case 0x08: gomp_notify(" FORCE_ALLOC "); break;
> - case 0x09: gomp_notify(" FORCE_TO "); break;
> - case 0x0a: gomp_notify(" FORCE_FROM "); break;
> - case 0x0b: gomp_notify(" FORCE_TOFROM "); break;
> - case 0x0c: gomp_notify(" FORCE_PRESENT "); break;
> - case 0x0d: gomp_notify(" FORCE_DEALLOC "); break;
> - case 0x0e: gomp_notify(" FORCE_DEVICEPTR "); break;
> + case 0x08: gomp_debug (0, " FORCE_ALLOC "); break;
> + case 0x09: gomp_debug (0, " FORCE_TO "); break;
> + case 0x0a: gomp_debug (0, " FORCE_FROM "); break;
> + case 0x0b: gomp_debug (0, " FORCE_TOFROM "); break;
> + case 0x0c: gomp_debug (0, " FORCE_PRESENT "); break;
> + case 0x0d: gomp_debug (0, " FORCE_DEALLOC "); break;
> + case 0x0e: gomp_debug (0, " FORCE_DEVICEPTR "); break;
>
> - case 0x18: gomp_notify(" FORCE_PRIVATE "); break;
> - case 0x19: gomp_notify(" FORCE_FIRSTPRIVATE "); break;
> -
> - case (unsigned char) -1: gomp_notify(" DUMMY "); break;
> - default: gomp_notify("UGH! 0x%x\n", kind);
> + case (unsigned char) -1: gomp_debug (0, " DUMMY "); break;
> + default: gomp_debug (0, "UGH! 0x%x\n", kind);
> }
>
> - gomp_notify("- %d - %4d/0x%04x ", 1 << (kind >> 8), (int)size, (int)size);
> - gomp_notify("- %p\n", hostaddr);
> -
> - return;
> + gomp_debug (0, "- %d - %4d/0x%04x ", 1 << (kind >> 8), (int) size, (int) size);
> + gomp_debug (0, "- %p\n", hostaddr);
> }
>
> static int
> @@ -92,7 +88,7 @@ find_pset (int pos, size_t mapnum, unsigned short *kinds)
> attribute_hidden void
> select_acc_device (int device_type)
> {
> - ACC_lazy_initialize ();
> + goacc_lazy_initialize ();
>
> if (device_type == GOMP_IF_CLAUSE_FALSE)
> return;
> @@ -109,7 +105,7 @@ select_acc_device (int device_type)
> }
> }
>
> -void goacc_wait (int async, int num_waits, va_list ap);
> +static void goacc_wait (int async, int num_waits, va_list ap);
>
> void
> GOACC_parallel (int device, void (*fn) (void *), const void *offload_table,
> @@ -133,8 +129,8 @@ GOACC_parallel (int device, void (*fn) (void *), const void *offload_table,
> gomp_fatal ("num_workers (%d) different from one is not yet supported",
> num_workers);
>
> - gomp_notify ("%s: mapnum=%zd, hostaddrs=%p, sizes=%p, kinds=%p, async=%d\n",
> - __FUNCTION__, mapnum, hostaddrs, sizes, kinds, async);
> + gomp_debug (0, "%s: mapnum=%zd, hostaddrs=%p, sizes=%p, kinds=%p, async=%d\n",
> + __FUNCTION__, mapnum, hostaddrs, sizes, kinds, async);
>
> select_acc_device (device);
>
> @@ -145,9 +141,9 @@ GOACC_parallel (int device, void (*fn) (void *), const void *offload_table,
> the host. */
> if (!if_clause_condition_value)
> {
> - ACC_save_and_set_bind (acc_device_host);
> + goacc_save_and_set_bind (acc_device_host);
> fn (hostaddrs);
> - ACC_restore_bind ();
> + goacc_restore_bind ();
> return;
> }
> else if (acc_device_type (acc_dev->type) == acc_device_host)
> @@ -213,8 +209,8 @@ GOACC_data_start (int device, const void *offload_table, size_t mapnum,
> bool if_clause_condition_value = device != GOMP_IF_CLAUSE_FALSE;
> struct target_mem_desc *tgt;
>
> - gomp_notify ("%s: mapnum=%zd, hostaddrs=%p, sizes=%p, kinds=%p\n",
> - __FUNCTION__, mapnum, hostaddrs, sizes, kinds);
> + gomp_debug (0, "%s: mapnum=%zd, hostaddrs=%p, sizes=%p, kinds=%p\n",
> + __FUNCTION__, mapnum, hostaddrs, sizes, kinds);
>
> select_acc_device (device);
>
> @@ -232,10 +228,10 @@ GOACC_data_start (int device, const void *offload_table, size_t mapnum,
> return;
> }
>
> - gomp_notify (" %s: prepare mappings\n", __FUNCTION__);
> + gomp_debug (0, " %s: prepare mappings\n", __FUNCTION__);
> tgt = gomp_map_vars (acc_dev, mapnum, hostaddrs, NULL, sizes, kinds, true,
> false);
> - gomp_notify (" %s: mappings prepared\n", __FUNCTION__);
> + gomp_debug (0, " %s: mappings prepared\n", __FUNCTION__);
> tgt->prev = thr->mapped_data;
> thr->mapped_data = tgt;
> }
> @@ -246,10 +242,10 @@ GOACC_data_end (void)
> struct goacc_thread *thr = goacc_thread ();
> struct target_mem_desc *tgt = thr->mapped_data;
>
> - gomp_notify (" %s: restore mappings\n", __FUNCTION__);
> + gomp_debug (0, " %s: restore mappings\n", __FUNCTION__);
> thr->mapped_data = tgt->prev;
> gomp_unmap_vars (tgt, true);
> - gomp_notify (" %s: mappings restored\n", __FUNCTION__);
> + gomp_debug (0, " %s: mappings restored\n", __FUNCTION__);
> }
>
> void
> @@ -397,8 +393,8 @@ GOACC_kernels (int device, void (*fn) (void *), const void *offload_table,
> int num_gangs, int num_workers, int vector_length,
> int async, int num_waits, ...)
> {
> - gomp_notify ("%s: mapnum=%zd, hostaddrs=%p, sizes=%p, kinds=%p\n", __FUNCTION__,
> - mapnum, hostaddrs, sizes, kinds);
> + gomp_debug (0, "%s: mapnum=%zd, hostaddrs=%p, sizes=%p, kinds=%p\n",
> + __FUNCTION__, mapnum, hostaddrs, sizes, kinds);
>
> va_list ap;
>
> @@ -411,12 +407,11 @@ GOACC_kernels (int device, void (*fn) (void *), const void *offload_table,
>
> va_end (ap);
>
> - /* TODO. */
> GOACC_parallel (device, fn, offload_table, mapnum, hostaddrs, sizes, kinds,
> - num_gangs, num_workers, vector_length, async, num_waits);
> + num_gangs, num_workers, vector_length, async, 0);
> }
>
> -void
> +static void
> goacc_wait (int async, int num_waits, va_list ap)
> {
> struct goacc_thread *thr = goacc_thread ();
> diff --git libgomp/oacc-plugin.c libgomp/oacc-plugin.c
> index 357cb5f..30ffe46 100644
> --- libgomp/oacc-plugin.c
> +++ libgomp/oacc-plugin.c
> @@ -2,7 +2,8 @@
>
> Contributed by Mentor Embedded.
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> diff --git libgomp/oacc-plugin.h libgomp/oacc-plugin.h
> index d05a28f..cad422a 100644
> --- libgomp/oacc-plugin.h
> +++ libgomp/oacc-plugin.h
> @@ -2,7 +2,8 @@
>
> Contributed by Mentor Embedded.
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> diff --git libgomp/omp.h.in libgomp/omp.h.in
> index 5caab5d..0683c37 100644
> --- libgomp/omp.h.in
> +++ libgomp/omp.h.in
> @@ -1,7 +1,8 @@
> /* Copyright (C) 2005-2014 Free Software Foundation, Inc.
> Contributed by Richard Henderson <rth@redhat.com>.
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> diff --git libgomp/omp_lib.f90.in libgomp/omp_lib.f90.in
> index 757053c..dab38e0 100644
> --- libgomp/omp_lib.f90.in
> +++ libgomp/omp_lib.f90.in
> @@ -1,7 +1,8 @@
> ! Copyright (C) 2005-2014 Free Software Foundation, Inc.
> ! Contributed by Jakub Jelinek <jakub@redhat.com>.
>
> -! This file is part of the GNU OpenMP Library (libgomp).
> +! 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
> diff --git libgomp/omp_lib.h.in libgomp/omp_lib.h.in
> index 691adb8..12c0f67 100644
> --- libgomp/omp_lib.h.in
> +++ libgomp/omp_lib.h.in
> @@ -1,7 +1,8 @@
> ! Copyright (C) 2005-2014 Free Software Foundation, Inc.
> ! Contributed by Jakub Jelinek <jakub@redhat.com>.
>
> -! This file is part of the GNU OpenMP Library (libgomp).
> +! 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
> diff --git libgomp/openacc.f90 libgomp/openacc.f90
> index e4d4d8f..c2952e7 100644
> --- libgomp/openacc.f90
> +++ libgomp/openacc.f90
> @@ -5,7 +5,8 @@
> ! Contributed by Tobias Burnus <burnus@net-b.de>
> ! and Mentor Embedded.
>
> -! This file is part of the GNU OpenMP Library (libgomp).
> +! 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
> diff --git libgomp/openacc.h libgomp/openacc.h
> index cf40d07..4ac1365 100644
> --- libgomp/openacc.h
> +++ libgomp/openacc.h
> @@ -2,9 +2,10 @@
>
> Copyright (C) 2013-2014 Free Software Foundation, Inc.
>
> - Contributed by Thomas Schwinge <thomas@codesourcery.com>.
> + Contributed by Mentor Embedded.
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> @@ -106,8 +107,8 @@ extern "C" {
> size_t, void **, size_t *, unsigned char *) __GOACC_NOTHROW;
> void ACC_add_device_code (void const *, char const *) __GOACC_NOTHROW;
>
> - void ACC_async_copy(int) __GOACC_NOTHROW;
> - void ACC_async_kern(int) __GOACC_NOTHROW;
> + void ACC_async_copy (int) __GOACC_NOTHROW;
> + void ACC_async_kern (int) __GOACC_NOTHROW;
>
> /* Old names. OpenACC does not specify whether these can or must
> not be macros, inlines or aliases for the new names. */
> diff --git libgomp/openacc_lib.h libgomp/openacc_lib.h
> index 4e335f2..1f630c9 100644
> --- libgomp/openacc_lib.h
> +++ libgomp/openacc_lib.h
> @@ -5,7 +5,8 @@
> ! Contributed by Tobias Burnus <burnus@net-b.de>
> ! and Mentor Embedded.
>
> -! This file is part of the GNU OpenMP Library (libgomp).
> +! 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
> diff --git libgomp/ordered.c libgomp/ordered.c
> index a9e1e20..39e5b2e 100644
> --- libgomp/ordered.c
> +++ libgomp/ordered.c
> @@ -1,7 +1,8 @@
> /* Copyright (C) 2005-2014 Free Software Foundation, Inc.
> Contributed by Richard Henderson <rth@redhat.com>.
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> diff --git libgomp/parallel.c libgomp/parallel.c
> index 40a1920..87dbc31 100644
> --- libgomp/parallel.c
> +++ libgomp/parallel.c
> @@ -1,7 +1,8 @@
> /* Copyright (C) 2005-2014 Free Software Foundation, Inc.
> Contributed by Richard Henderson <rth@redhat.com>.
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> diff --git libgomp/plugin/Makefile.in libgomp/plugin/Makefile.in
> deleted file mode 100644
> index 1e5cb9d..0000000
> --- libgomp/plugin/Makefile.in
> +++ /dev/null
> [...]
> diff --git libgomp/plugin/Makefile.am libgomp/plugin/Makefrag.am
> similarity index 70%
> rename from libgomp/plugin/Makefile.am
> rename to libgomp/plugin/Makefrag.am
> index 59a5b95..7efea3a 100644
> --- libgomp/plugin/Makefile.am
> +++ libgomp/plugin/Makefrag.am
> @@ -1,10 +1,11 @@
> -# Plugins for offload execution.
> +# Plugins for offload execution, Makefile.am fragment.
> #
> # Copyright (C) 2014 Free Software Foundation, Inc.
> #
> # Contributed by Mentor Embedded.
> #
> -# This file is part of the GNU OpenMP Library (libgomp).
> +# 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
> @@ -25,23 +26,11 @@
> # see the files COPYING3 and COPYING.RUNTIME respectively. If not, see
> # <http://www.gnu.org/licenses/>.
>
> -ACLOCAL_AMFLAGS = -I ../.. -I ../../config
> -
> -config_path = @config_path@
> -search_path = .. $(addprefix $(top_srcdir)/../config/, $(config_path)) \
> - $(top_srcdir) $(top_srcdir)/../../include $(top_srcdir)/..
> -
> -AM_CPPFLAGS = $(addprefix -I, $(search_path))
> -AM_CFLAGS = $(XCFLAGS)
> -AM_LDFLAGS = $(XLDFLAGS) $(SECTION_LDFLAGS) $(OPT_LDFLAGS)
> -
> -toolexeclib_LTLIBRARIES =
> -
> if PLUGIN_NVPTX
> # Nvidia PTX OpenACC plugin.
> libgomp_plugin_nvptx_version_info = -version-info $(libtool_VERSION)
> toolexeclib_LTLIBRARIES += libgomp-plugin-nvptx.la
> -libgomp_plugin_nvptx_la_SOURCES = plugin-nvptx.c
> +libgomp_plugin_nvptx_la_SOURCES = plugin/plugin-nvptx.c
> libgomp_plugin_nvptx_la_CPPFLAGS = $(AM_CPPFLAGS) $(PLUGIN_NVPTX_CPPFLAGS)
> libgomp_plugin_nvptx_la_LDFLAGS = $(libgomp_plugin_nvptx_version_info) \
> $(lt_host_flags)
> @@ -52,13 +41,8 @@ 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-host.c
> +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_LIBTOOLFLAGS = --tag=disable-static
> -
> -LTLDFLAGS = $(shell $(SHELL) $(top_srcdir)/../../libtool-ldflags $(LDFLAGS))
> -
> -LINK = $(LIBTOOL) --tag CC $(AM_LIBTOOLFLAGS) $(LIBTOOLFLAGS) --mode=link \
> - $(CCLD) $(AM_CFLAGS) $(CFLAGS) $(AM_LDFLAGS) $(LTLDFLAGS) -o $@
> diff --git libgomp/plugin/aclocal.m4 libgomp/plugin/aclocal.m4
> deleted file mode 100644
> index 06820b7..0000000
> --- libgomp/plugin/aclocal.m4
> +++ /dev/null
> [...]
> diff --git libgomp/plugin/config.h.in libgomp/plugin/config.h.in
> deleted file mode 100644
> index d044b92..0000000
> --- libgomp/plugin/config.h.in
> +++ /dev/null
> [...]
> diff --git libgomp/plugin/configfrag.ac libgomp/plugin/configfrag.ac
> new file mode 100644
> index 0000000..06c2be6
> --- /dev/null
> +++ libgomp/plugin/configfrag.ac
> @@ -0,0 +1,148 @@
> +# Plugins for offload execution, configure.ac fragment.
> +#
> +# Copyright (C) 2014 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/>.
> +
> +offload_targets=
> +AC_SUBST(offload_targets)
> +plugin_support=yes
> +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
> +
> +# Look for the CUDA driver package.
> +CUDA_DRIVER_INCLUDE=
> +CUDA_DRIVER_LIB=
> +AC_SUBST(CUDA_DRIVER_INCLUDE)
> +AC_SUBST(CUDA_DRIVER_LIB)
> +CUDA_DRIVER_CPPFLAGS=
> +CUDA_DRIVER_LDFLAGS=
> +AC_ARG_WITH(cuda-driver,
> + [AS_HELP_STRING([--with-cuda-driver=PATH],
> + [specify prefix directory for installed CUDA driver package.
> + Equivalent to --with-cuda-driver-include=PATH/include
> + plus --with-cuda-driver-lib=PATH/lib])])
> +AC_ARG_WITH(cuda-driver-include,
> + [AS_HELP_STRING([--with-cuda-driver-include=PATH],
> + [specify directory for installed CUDA driver include files])])
> +AC_ARG_WITH(cuda-driver-lib,
> + [AS_HELP_STRING([--with-cuda-driver-lib=PATH],
> + [specify directory for the installed CUDA driver library])])
> +if test "x$with_cuda_driver" != x; then
> + CUDA_DRIVER_INCLUDE=$with_cuda_driver/include
> + CUDA_DRIVER_LIB=$with_cuda_driver/lib
> +fi
> +if test "x$with_cuda_driver_include" != x; then
> + CUDA_DRIVER_INCLUDE=$with_cuda_driver_include
> +fi
> +if test "x$with_cuda_driver_lib" != x; then
> + CUDA_DRIVER_LIB=$with_cuda_driver_lib
> +fi
> +if test "x$CUDA_DRIVER_INCLUDE" != x; then
> + CUDA_DRIVER_CPPFLAGS=-I$CUDA_DRIVER_INCLUDE
> +fi
> +if test "x$CUDA_DRIVER_LIB" != x; then
> + CUDA_DRIVER_LDFLAGS=-L$CUDA_DRIVER_LIB
> +fi
> +
> +PLUGIN_NVPTX=0
> +PLUGIN_NVPTX_CPPFLAGS=
> +PLUGIN_NVPTX_LDFLAGS=
> +PLUGIN_NVPTX_LIBS=
> +AC_SUBST(PLUGIN_NVPTX)
> +AC_SUBST(PLUGIN_NVPTX_CPPFLAGS)
> +AC_SUBST(PLUGIN_NVPTX_LDFLAGS)
> +AC_SUBST(PLUGIN_NVPTX_LIBS)
> +
> +# Get offload targets and path to install tree of offloading compiler.
> +offload_additional_options=
> +offload_additional_lib_paths=
> +AC_SUBST(offload_additional_options)
> +AC_SUBST(offload_additional_lib_paths)
> +if test x"$enable_offload_targets" != x; then
> + for tgt in `echo $enable_offload_targets | sed -e 's#,# #g'`; do
> + tgt_dir=`echo $tgt | grep '=' | sed 's/.*=//'`
> + tgt=`echo $tgt | sed 's/=.*//'`
> + case $tgt in
> + *-intelmic-* | *-intelmicemul-*)
> + tgt_name=intelmic
> + ;;
> + nvptx*)
> + tgt_name=nvptx
> + PLUGIN_NVPTX=$tgt
> + PLUGIN_NVPTX_CPPFLAGS=$CUDA_DRIVER_CPPFLAGS
> + PLUGIN_NVPTX_LDFLAGS=$CUDA_DRIVER_LDFLAGS
> + PLUGIN_NVPTX_LIBS='-lcuda'
> +
> + PLUGIN_NVPTX_save_CPPFLAGS=$CPPFLAGS
> + CPPFLAGS="$PLUGIN_NVPTX_CPPFLAGS $CPPFLAGS"
> + PLUGIN_NVPTX_save_LDFLAGS=$LDFLAGS
> + LDFLAGS="$PLUGIN_NVPTX_LDFLAGS $LDFLAGS"
> + PLUGIN_NVPTX_save_LIBS=$LIBS
> + LIBS="$PLUGIN_NVPTX_LIBS $LIBS"
> + AC_LINK_IFELSE(
> + [AC_LANG_PROGRAM(
> + [#include "cuda.h"],
> + [CUresult r = cuCtxPushCurrent (NULL);])],
> + [PLUGIN_NVPTX=1])
> + CPPFLAGS=$PLUGIN_NVPTX_save_CPPFLAGS
> + LDFLAGS=$PLUGIN_NVPTX_save_LDFLAGS
> + LIBS=$PLUGIN_NVPTX_save_LIBS
> + case $PLUGIN_NVPTX in
> + nvptx*)
> + PLUGIN_NVPTX=0
> + AC_MSG_ERROR([CUDA driver package required for nvptx support])
> + ;;
> + esac
> + ;;
> + *)
> + AC_MSG_ERROR([unknown offload target specified])
> + ;;
> + esac
> + if test x"$offload_targets" = x; then
> + offload_targets=$tgt_name
> + else
> + offload_targets=$offload_targets,$tgt_name
> + fi
> + if test x"$tgt_dir" != x; then
> + offload_additional_options="$offload_additional_options -B$tgt_dir/libexec/gcc/\$(target_alias)/\$(gcc_version) -B$tgt_dir/bin"
> + offload_additional_lib_paths="$offload_additional_lib_paths:$tgt_dir/lib64:$tgt_dir/lib"
> + else
> + offload_additional_options="$offload_additional_options -B\$(libexecdir)/gcc/\$(target_alias)/\$(gcc_version) -B\$(bindir)"
> + offload_additional_lib_paths="$offload_additional_lib_paths:$toolexeclibdir"
> + fi
> + done
> +fi
> +AC_DEFINE_UNQUOTED(OFFLOAD_TARGETS, "$offload_targets",
> + [Define to hold the list of target names suitable for offloading.])
> +AM_CONDITIONAL([PLUGIN_NVPTX], [test $PLUGIN_NVPTX = 1])
> +AC_DEFINE_UNQUOTED([PLUGIN_NVPTX], [$PLUGIN_NVPTX],
> + [Define to 1 if the NVIDIA plugin is built, 0 if not.])
> diff --git libgomp/plugin/configure libgomp/plugin/configure
> deleted file mode 100644
> index f43bc38..0000000
> --- libgomp/plugin/configure
> +++ /dev/null
> [...]
> diff --git libgomp/plugin/configure.ac libgomp/plugin/configure.ac
> deleted file mode 100644
> index bc2565c..0000000
> --- libgomp/plugin/configure.ac
> +++ /dev/null
> [...]
> diff --git libgomp/plugin/plugin-host.c libgomp/plugin/plugin-host.c
> index f92aaee..8bca998 100644
> --- libgomp/plugin/plugin-host.c
> +++ libgomp/plugin/plugin-host.c
> @@ -4,7 +4,8 @@
>
> Contributed by Mentor Embedded.
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> @@ -62,10 +63,6 @@ static struct gomp_device_descr host_dispatch;
> STATIC const char *
> GOMP_OFFLOAD_get_name (void)
> {
> -#ifdef DEBUG
> - fprintf (stderr, SELF "%s:%s\n", __FILE__, __FUNCTION__);
> -#endif
> -
> #ifdef HOST_NONSHM_PLUGIN
> return "host_nonshm";
> #else
> @@ -76,10 +73,6 @@ GOMP_OFFLOAD_get_name (void)
> STATIC int
> GOMP_OFFLOAD_get_type (void)
> {
> -#ifdef DEBUG
> - fprintf (stderr, SELF "%s:%s\n", __FILE__, __FUNCTION__);
> -#endif
> -
> #ifdef HOST_NONSHM_PLUGIN
> return OFFLOAD_TARGET_TYPE_HOST_NONSHM;
> #else
> @@ -90,27 +83,18 @@ GOMP_OFFLOAD_get_type (void)
> STATIC unsigned int
> GOMP_OFFLOAD_get_caps (void)
> {
> - unsigned int caps = TARGET_CAP_OPENACC_200 | TARGET_CAP_OPENMP_400
> - | TARGET_CAP_NATIVE_EXEC;
> + unsigned int caps = TARGET_CAP_OPENACC_200 | TARGET_CAP_NATIVE_EXEC;
>
> #ifndef HOST_NONSHM_PLUGIN
> caps |= TARGET_CAP_SHARED_MEM;
> #endif
>
> -#ifdef DEBUG
> - fprintf (stderr, SELF "%s:%s: 0x%x\n", __FILE__, __FUNCTION__, caps);
> -#endif
> -
> return caps;
> }
>
> STATIC int
> GOMP_OFFLOAD_get_num_devices (void)
> {
> -#ifdef DEBUG
> - fprintf (stderr, SELF "%s:%s\n", __FILE__, __FUNCTION__);
> -#endif
> -
> return 1;
> }
>
> @@ -118,76 +102,46 @@ STATIC void
> GOMP_OFFLOAD_register_image (void *host_table __attribute__((unused)),
> void *target_data __attribute__((unused)))
> {
> -#ifdef DEBUG
> - fprintf (stderr, SELF "%s:%s (%p, %p)\n", __FILE__, __FUNCTION__, host_table,
> - target_data);
> -#endif
> }
>
> STATIC void
> GOMP_OFFLOAD_init_device (int n __attribute__((unused)))
> {
> -#ifdef DEBUG
> - fprintf (stderr, SELF "%s:%s\n", __FILE__, __FUNCTION__);
> -#endif
> }
>
> STATIC void
> GOMP_OFFLOAD_fini_device (int n __attribute__((unused)))
> {
> -#ifdef DEBUG
> - fprintf (stderr, SELF "%s:%s\n", __FILE__, __FUNCTION__);
> -#endif
> }
>
> STATIC int
> GOMP_OFFLOAD_get_table (int n __attribute__((unused)),
> struct mapping_table **table __attribute__((unused)))
> {
> -#ifdef DEBUG
> - fprintf (stderr, SELF "%s:%s (%p)\n", __FILE__, __FUNCTION__, table);
> -#endif
> -
> return 0;
> }
>
> STATIC void *
> GOMP_OFFLOAD_openacc_open_device (int n)
> {
> -#ifdef DEBUG
> - fprintf (stderr, SELF "%s:%s (%u)\n", __FILE__, __FUNCTION__, n);
> -#endif
> -
> return (void *) (intptr_t) n;
> }
>
> STATIC int
> GOMP_OFFLOAD_openacc_close_device (void *hnd)
> {
> -#ifdef DEBUG
> - fprintf (stderr, SELF "%s:%s (%p)\n", __FILE__, __FUNCTION__, hnd);
> -#endif
> -
> return 0;
> }
>
> STATIC int
> GOMP_OFFLOAD_openacc_get_device_num (void)
> {
> -#ifdef DEBUG
> - fprintf (stderr, SELF "%s:%s\n", __FILE__, __FUNCTION__);
> -#endif
> -
> return 0;
> }
>
> STATIC void
> GOMP_OFFLOAD_openacc_set_device_num (int n)
> {
> -#ifdef DEBUG
> - fprintf (stderr, SELF "%s:%s (%u)\n", __FILE__, __FUNCTION__, n);
> -#endif
> -
> if (n > 0)
> GOMP(fatal) ("device number %u out of range for host execution", n);
> }
> @@ -195,22 +149,12 @@ GOMP_OFFLOAD_openacc_set_device_num (int n)
> STATIC void *
> GOMP_OFFLOAD_alloc (int n __attribute__((unused)), size_t s)
> {
> - void *ptr = GOMP(malloc) (s);
> -
> -#ifdef DEBUG
> - fprintf (stderr, SELF "%s:%s (%zd): %p\n", __FILE__, __FUNCTION__, s, ptr);
> -#endif
> -
> - return ptr;
> + return GOMP(malloc) (s);
> }
>
> STATIC void
> GOMP_OFFLOAD_free (int n __attribute__((unused)), void *p)
> {
> -#ifdef DEBUG
> - fprintf (stderr, SELF "%s:%s (%p)\n", __FILE__, __FUNCTION__, p);
> -#endif
> -
> free (p);
> }
>
> @@ -218,11 +162,6 @@ STATIC void *
> GOMP_OFFLOAD_host2dev (int n __attribute__((unused)), void *d, const void *h,
> size_t s)
> {
> -#ifdef DEBUG
> - fprintf (stderr, SELF "%s:%s (%p, %p, %zd)\n", __FILE__, __FUNCTION__, d, h,
> - s);
> -#endif
> -
> #ifdef HOST_NONSHM_PLUGIN
> memcpy (d, h, s);
> #endif
> @@ -234,11 +173,6 @@ STATIC void *
> GOMP_OFFLOAD_dev2host (int n __attribute__((unused)), void *h, const void *d,
> size_t s)
> {
> -#ifdef DEBUG
> - fprintf (stderr, SELF "%s:%s (%p, %p, %zd)\n", __FILE__, __FUNCTION__, h, d,
> - s);
> -#endif
> -
> #ifdef HOST_NONSHM_PLUGIN
> memcpy (h, d, s);
> #endif
> @@ -249,11 +183,6 @@ GOMP_OFFLOAD_dev2host (int n __attribute__((unused)), void *h, const void *d,
> STATIC void
> GOMP_OFFLOAD_run (int n __attribute__((unused)), void *fn_ptr, void *vars)
> {
> -#ifdef DEBUG
> - fprintf (stderr, SELF "%s:%s (%p, %p)\n", __FILE__, __FUNCTION__, fn_ptr,
> - vars);
> -#endif
> -
> void (*fn)(void *) = (void (*)(void *)) fn_ptr;
>
> fn (vars);
> @@ -272,12 +201,6 @@ GOMP_OFFLOAD_openacc_parallel (void (*fn) (void *),
> int async __attribute__((unused)),
> void *targ_mem_desc __attribute__((unused)))
> {
> -#ifdef DEBUG
> - fprintf (stderr, SELF "%s:%s (%p, %zu, %p, %p, %p, %d, %d, %d, %d, %p)\n",
> - __FILE__, __FUNCTION__, fn, mapnum, hostaddrs, sizes, kinds,
> - num_gangs, num_workers, vector_length, async, targ_mem_desc);
> -#endif
> -
> #ifdef HOST_NONSHM_PLUGIN
> fn (devaddrs);
> #else
> @@ -298,63 +221,39 @@ GOMP_OFFLOAD_openacc_register_async_cleanup (void *targ_mem_desc)
> STATIC void
> GOMP_OFFLOAD_openacc_async_set_async (int async __attribute__((unused)))
> {
> -#ifdef DEBUG
> - fprintf (stderr, SELF "%s:%s (%d)\n", __FILE__, __FUNCTION__, async);
> -#endif
> }
>
> STATIC int
> GOMP_OFFLOAD_openacc_async_test (int async __attribute__((unused)))
> {
> -#ifdef DEBUG
> - fprintf (stderr, SELF "%s:%s (%d)\n", __FILE__, __FUNCTION__, async);
> -#endif
> -
> return 1;
> }
>
> STATIC int
> GOMP_OFFLOAD_openacc_async_test_all (void)
> {
> -#ifdef DEBUG
> - fprintf (stderr, SELF "%s:%s\n", __FILE__, __FUNCTION__);
> -#endif
> -
> return 1;
> }
>
> STATIC void
> GOMP_OFFLOAD_openacc_async_wait (int async __attribute__((unused)))
> {
> -#ifdef DEBUG
> - fprintf (stderr, SELF "%s:%s (%d)\n", __FILE__, __FUNCTION__, async);
> -#endif
> }
>
> STATIC void
> GOMP_OFFLOAD_openacc_async_wait_all (void)
> {
> -#ifdef DEBUG
> - fprintf (stderr, SELF "%s:%s\n", __FILE__, __FUNCTION__);
> -#endif
> }
>
> STATIC void
> GOMP_OFFLOAD_openacc_async_wait_async (int async1 __attribute__((unused)),
> int async2 __attribute__((unused)))
> {
> -#ifdef DEBUG
> - fprintf (stderr, SELF "%s:%s (%d, %d)\n", __FILE__, __FUNCTION__, async1,
> - async2);
> -#endif
> }
>
> STATIC void
> GOMP_OFFLOAD_openacc_async_wait_all_async (int async __attribute__((unused)))
> {
> -#ifdef DEBUG
> - fprintf (stderr, SELF "%s:%s (%d)\n", __FILE__, __FUNCTION__, async);
> -#endif
> }
>
> STATIC void *
> diff --git libgomp/plugin/plugin-nvptx.c libgomp/plugin/plugin-nvptx.c
> index 320ba6b5..bc5739a 100644
> --- libgomp/plugin/plugin-nvptx.c
> +++ libgomp/plugin/plugin-nvptx.c
> @@ -4,7 +4,8 @@
>
> Contributed by Mentor Embedded.
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> @@ -48,76 +49,77 @@
>
> #define ARRAYSIZE(X) (sizeof (X) / sizeof ((X)[0]))
>
> -static struct _errlist
> +static struct
> {
> CUresult r;
> char *m;
> -} cuErrorList[] = {
> - { CUDA_ERROR_INVALID_VALUE, "invalid value" },
> - { CUDA_ERROR_OUT_OF_MEMORY, "out of memory" },
> - { CUDA_ERROR_NOT_INITIALIZED, "not initialized" },
> - { CUDA_ERROR_DEINITIALIZED, "deinitialized" },
> - { CUDA_ERROR_PROFILER_DISABLED, "profiler disabled" },
> - { CUDA_ERROR_PROFILER_NOT_INITIALIZED, "profiler not initialized" },
> - { CUDA_ERROR_PROFILER_ALREADY_STARTED, "already started" },
> - { CUDA_ERROR_PROFILER_ALREADY_STOPPED, "already stopped" },
> - { CUDA_ERROR_NO_DEVICE, "no device" },
> - { CUDA_ERROR_INVALID_DEVICE, "invalid device" },
> - { CUDA_ERROR_INVALID_IMAGE, "invalid image" },
> - { CUDA_ERROR_INVALID_CONTEXT, "invalid context" },
> - { CUDA_ERROR_CONTEXT_ALREADY_CURRENT, "context already current" },
> - { CUDA_ERROR_MAP_FAILED, "map error" },
> - { CUDA_ERROR_UNMAP_FAILED, "unmap error" },
> - { CUDA_ERROR_ARRAY_IS_MAPPED, "array is mapped" },
> - { CUDA_ERROR_ALREADY_MAPPED, "already mapped" },
> - { CUDA_ERROR_NO_BINARY_FOR_GPU, "no binary for gpu" },
> - { CUDA_ERROR_ALREADY_ACQUIRED, "already acquired" },
> - { CUDA_ERROR_NOT_MAPPED, "not mapped" },
> - { CUDA_ERROR_NOT_MAPPED_AS_ARRAY, "not mapped as array" },
> - { CUDA_ERROR_NOT_MAPPED_AS_POINTER, "not mapped as pointer" },
> - { CUDA_ERROR_ECC_UNCORRECTABLE, "ecc uncorrectable" },
> - { CUDA_ERROR_UNSUPPORTED_LIMIT, "unsupported limit" },
> - { CUDA_ERROR_CONTEXT_ALREADY_IN_USE, "context already in use" },
> - { CUDA_ERROR_PEER_ACCESS_UNSUPPORTED, "peer access unsupported" },
> - { CUDA_ERROR_INVALID_SOURCE, "invalid source" },
> - { CUDA_ERROR_FILE_NOT_FOUND, "file not found" },
> - { CUDA_ERROR_SHARED_OBJECT_SYMBOL_NOT_FOUND,
> - "shared object symbol not found" },
> - { CUDA_ERROR_SHARED_OBJECT_INIT_FAILED, "shared object init error" },
> - { CUDA_ERROR_OPERATING_SYSTEM, "operating system" },
> - { CUDA_ERROR_INVALID_HANDLE, "invalid handle" },
> - { CUDA_ERROR_NOT_FOUND, "not found" },
> - { CUDA_ERROR_NOT_READY, "not ready" },
> - { CUDA_ERROR_LAUNCH_FAILED, "launch error" },
> - { CUDA_ERROR_LAUNCH_OUT_OF_RESOURCES, "launch out of resources" },
> - { CUDA_ERROR_LAUNCH_TIMEOUT, "launch timeout" },
> - { CUDA_ERROR_LAUNCH_INCOMPATIBLE_TEXTURING,
> - "launch incompatibe texturing" },
> - { CUDA_ERROR_PEER_ACCESS_ALREADY_ENABLED, "peer access already enabled" },
> - { CUDA_ERROR_PEER_ACCESS_NOT_ENABLED, "peer access not enabled " },
> - { CUDA_ERROR_PRIMARY_CONTEXT_ACTIVE, "primary cotext active" },
> - { CUDA_ERROR_CONTEXT_IS_DESTROYED, "context is destroyed" },
> - { CUDA_ERROR_ASSERT, "assert" },
> - { CUDA_ERROR_TOO_MANY_PEERS, "too many peers" },
> - { CUDA_ERROR_HOST_MEMORY_ALREADY_REGISTERED,
> - "host memory already registered" },
> - { CUDA_ERROR_HOST_MEMORY_NOT_REGISTERED, "host memory not registered" },
> - { CUDA_ERROR_NOT_PERMITTED, "no permitted" },
> - { CUDA_ERROR_NOT_SUPPORTED, "not supported" },
> - { CUDA_ERROR_UNKNOWN, "unknown" }
> +} cuda_errlist[]=
> +{
> + { CUDA_ERROR_INVALID_VALUE, "invalid value" },
> + { CUDA_ERROR_OUT_OF_MEMORY, "out of memory" },
> + { CUDA_ERROR_NOT_INITIALIZED, "not initialized" },
> + { CUDA_ERROR_DEINITIALIZED, "deinitialized" },
> + { CUDA_ERROR_PROFILER_DISABLED, "profiler disabled" },
> + { CUDA_ERROR_PROFILER_NOT_INITIALIZED, "profiler not initialized" },
> + { CUDA_ERROR_PROFILER_ALREADY_STARTED, "already started" },
> + { CUDA_ERROR_PROFILER_ALREADY_STOPPED, "already stopped" },
> + { CUDA_ERROR_NO_DEVICE, "no device" },
> + { CUDA_ERROR_INVALID_DEVICE, "invalid device" },
> + { CUDA_ERROR_INVALID_IMAGE, "invalid image" },
> + { CUDA_ERROR_INVALID_CONTEXT, "invalid context" },
> + { CUDA_ERROR_CONTEXT_ALREADY_CURRENT, "context already current" },
> + { CUDA_ERROR_MAP_FAILED, "map error" },
> + { CUDA_ERROR_UNMAP_FAILED, "unmap error" },
> + { CUDA_ERROR_ARRAY_IS_MAPPED, "array is mapped" },
> + { CUDA_ERROR_ALREADY_MAPPED, "already mapped" },
> + { CUDA_ERROR_NO_BINARY_FOR_GPU, "no binary for gpu" },
> + { CUDA_ERROR_ALREADY_ACQUIRED, "already acquired" },
> + { CUDA_ERROR_NOT_MAPPED, "not mapped" },
> + { CUDA_ERROR_NOT_MAPPED_AS_ARRAY, "not mapped as array" },
> + { CUDA_ERROR_NOT_MAPPED_AS_POINTER, "not mapped as pointer" },
> + { CUDA_ERROR_ECC_UNCORRECTABLE, "ecc uncorrectable" },
> + { CUDA_ERROR_UNSUPPORTED_LIMIT, "unsupported limit" },
> + { CUDA_ERROR_CONTEXT_ALREADY_IN_USE, "context already in use" },
> + { CUDA_ERROR_PEER_ACCESS_UNSUPPORTED, "peer access unsupported" },
> + { CUDA_ERROR_INVALID_SOURCE, "invalid source" },
> + { CUDA_ERROR_FILE_NOT_FOUND, "file not found" },
> + { CUDA_ERROR_SHARED_OBJECT_SYMBOL_NOT_FOUND,
> + "shared object symbol not found" },
> + { CUDA_ERROR_SHARED_OBJECT_INIT_FAILED, "shared object init error" },
> + { CUDA_ERROR_OPERATING_SYSTEM, "operating system" },
> + { CUDA_ERROR_INVALID_HANDLE, "invalid handle" },
> + { CUDA_ERROR_NOT_FOUND, "not found" },
> + { CUDA_ERROR_NOT_READY, "not ready" },
> + { CUDA_ERROR_LAUNCH_FAILED, "launch error" },
> + { CUDA_ERROR_LAUNCH_OUT_OF_RESOURCES, "launch out of resources" },
> + { CUDA_ERROR_LAUNCH_TIMEOUT, "launch timeout" },
> + { CUDA_ERROR_LAUNCH_INCOMPATIBLE_TEXTURING,
> + "launch incompatibe texturing" },
> + { CUDA_ERROR_PEER_ACCESS_ALREADY_ENABLED, "peer access already enabled" },
> + { CUDA_ERROR_PEER_ACCESS_NOT_ENABLED, "peer access not enabled " },
> + { CUDA_ERROR_PRIMARY_CONTEXT_ACTIVE, "primary cotext active" },
> + { CUDA_ERROR_CONTEXT_IS_DESTROYED, "context is destroyed" },
> + { CUDA_ERROR_ASSERT, "assert" },
> + { CUDA_ERROR_TOO_MANY_PEERS, "too many peers" },
> + { CUDA_ERROR_HOST_MEMORY_ALREADY_REGISTERED,
> + "host memory already registered" },
> + { CUDA_ERROR_HOST_MEMORY_NOT_REGISTERED, "host memory not registered" },
> + { CUDA_ERROR_NOT_PERMITTED, "not permitted" },
> + { CUDA_ERROR_NOT_SUPPORTED, "not supported" },
> + { CUDA_ERROR_UNKNOWN, "unknown" }
> };
>
> static char errmsg[128];
>
> static char *
> -cuErrorMsg (CUresult r)
> +cuda_error (CUresult r)
> {
> int i;
>
> - for (i = 0; i < ARRAYSIZE (cuErrorList); i++)
> + for (i = 0; i < ARRAYSIZE (cuda_errlist); i++)
> {
> - if (cuErrorList[i].r == r)
> - return &cuErrorList[i].m[0];
> + if (cuda_errlist[i].r == r)
> + return &cuda_errlist[i].m[0];
> }
>
> sprintf (&errmsg[0], "unknown result code: %5d", r);
> @@ -131,9 +133,9 @@ struct targ_fn_descriptor
> const char *name;
> };
>
> -static bool PTX_inited = false;
> +static bool ptx_inited = false;
>
> -struct PTX_stream
> +struct ptx_stream
> {
> CUstream stream;
> pthread_t host_thread;
> @@ -147,15 +149,15 @@ struct PTX_stream
> void *h_prev;
> void *h_tail;
>
> - struct PTX_stream *next;
> + struct ptx_stream *next;
> };
>
> /* Thread-specific data for PTX. */
>
> struct nvptx_thread
> {
> - struct PTX_stream *current_stream;
> - struct PTX_device *ptx_dev;
> + struct ptx_stream *current_stream;
> + struct ptx_device *ptx_dev;
> };
>
> struct map
> @@ -166,7 +168,7 @@ struct map
> };
>
> static void
> -map_init (struct PTX_stream *s)
> +map_init (struct ptx_stream *s)
> {
> CUresult r;
>
> @@ -178,11 +180,11 @@ map_init (struct PTX_stream *s)
>
> r = cuMemAllocHost (&s->h, size);
> if (r != CUDA_SUCCESS)
> - GOMP_PLUGIN_fatal ("cuMemAllocHost error: %s", cuErrorMsg (r));
> + GOMP_PLUGIN_fatal ("cuMemAllocHost error: %s", cuda_error (r));
>
> r = cuMemHostGetDevicePointer (&s->d, s->h, 0);
> if (r != CUDA_SUCCESS)
> - GOMP_PLUGIN_fatal ("cuMemHostGetDevicePointer error: %s", cuErrorMsg (r));
> + GOMP_PLUGIN_fatal ("cuMemHostGetDevicePointer error: %s", cuda_error (r));
>
> assert (s->h);
>
> @@ -195,17 +197,17 @@ map_init (struct PTX_stream *s)
> }
>
> static void
> -map_fini (struct PTX_stream *s)
> +map_fini (struct ptx_stream *s)
> {
> CUresult r;
> -
> +
> r = cuMemFreeHost (s->h);
> if (r != CUDA_SUCCESS)
> - GOMP_PLUGIN_fatal ("cuMemFreeHost error: %s", cuErrorMsg (r));
> + GOMP_PLUGIN_fatal ("cuMemFreeHost error: %s", cuda_error (r));
> }
>
> static void
> -map_pop (struct PTX_stream *s)
> +map_pop (struct ptx_stream *s)
> {
> struct map *m;
>
> @@ -234,7 +236,7 @@ map_pop (struct PTX_stream *s)
> }
>
> static void
> -map_push (struct PTX_stream *s, int async, size_t size, void **h, void **d)
> +map_push (struct ptx_stream *s, int async, size_t size, void **h, void **d)
> {
> int left;
> int offset;
> @@ -285,18 +287,18 @@ map_push (struct PTX_stream *s, int async, size_t size, void **h, void **d)
> return;
> }
>
> -struct PTX_device
> +struct ptx_device
> {
> CUcontext ctx;
> bool ctx_shared;
> CUdevice dev;
> - struct PTX_stream *null_stream;
> + struct ptx_stream *null_stream;
> /* All non-null streams associated with this device (actually context),
> either created implicitly or passed in from the user (via
> acc_set_cuda_stream). */
> - struct PTX_stream *active_streams;
> + struct ptx_stream *active_streams;
> struct {
> - struct PTX_stream **arr;
> + struct ptx_stream **arr;
> int size;
> } async_streams;
> /* A lock for use when manipulating the above stream list and array. */
> @@ -308,10 +310,10 @@ struct PTX_device
> int mode;
> bool mkern;
>
> - struct PTX_device *next;
> + struct ptx_device *next;
> };
>
> -enum PTX_event_type
> +enum ptx_event_type
> {
> PTX_EVT_MEM,
> PTX_EVT_KNL,
> @@ -319,18 +321,18 @@ enum PTX_event_type
> PTX_EVT_ASYNC_CLEANUP
> };
>
> -struct PTX_event
> +struct ptx_event
> {
> CUevent *evt;
> int type;
> void *addr;
> int ord;
>
> - struct PTX_event *next;
> + struct ptx_event *next;
> };
>
> -static gomp_mutex_t PTX_event_lock;
> -static struct PTX_event *PTX_events;
> +static gomp_mutex_t ptx_event_lock;
> +static struct ptx_event *ptx_events;
>
> #define _XSTR(s) _STR(s)
> #define _STR(s) #s
> @@ -338,44 +340,44 @@ static struct PTX_event *PTX_events;
> static struct _synames
> {
> char *n;
> -} cuSymNames[] =
> +} 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) }
> + { _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
> @@ -388,15 +390,15 @@ verify_device_library (void)
> if (!dh)
> return -1;
>
> - for (i = 0; i < ARRAYSIZE (cuSymNames); i++)
> + for (i = 0; i < ARRAYSIZE (cuda_symnames); i++)
> {
> - ds = dlsym (dh, cuSymNames[i].n);
> + ds = dlsym (dh, cuda_symnames[i].n);
> if (!ds)
> return -1;
> }
>
> dlclose (dh);
> -
> +
> return 0;
> }
>
> @@ -407,11 +409,11 @@ nvptx_thread (void)
> }
>
> static void
> -init_streams_for_device (struct PTX_device *ptx_dev, int concurrency)
> +init_streams_for_device (struct ptx_device *ptx_dev, int concurrency)
> {
> int i;
> - struct PTX_stream *null_stream
> - = GOMP_PLUGIN_malloc (sizeof (struct PTX_stream));
> + struct ptx_stream *null_stream
> + = GOMP_PLUGIN_malloc (sizeof (struct ptx_stream));
>
> null_stream->stream = NULL;
> null_stream->host_thread = pthread_self ();
> @@ -420,39 +422,39 @@ init_streams_for_device (struct PTX_device *ptx_dev, int concurrency)
> null_stream->h = NULL;
> map_init (null_stream);
> ptx_dev->null_stream = null_stream;
> -
> +
> ptx_dev->active_streams = NULL;
> GOMP_PLUGIN_mutex_init (&ptx_dev->stream_lock);
> -
> +
> if (concurrency < 1)
> concurrency = 1;
> -
> +
> /* This is just a guess -- make space for as many async streams as the
> current device is capable of concurrently executing. This can grow
> later as necessary. No streams are created yet. */
> ptx_dev->async_streams.arr
> - = GOMP_PLUGIN_malloc (concurrency * sizeof (struct PTX_stream *));
> + = GOMP_PLUGIN_malloc (concurrency * sizeof (struct ptx_stream *));
> ptx_dev->async_streams.size = concurrency;
> -
> +
> for (i = 0; i < concurrency; i++)
> ptx_dev->async_streams.arr[i] = NULL;
> }
>
> static void
> -fini_streams_for_device (struct PTX_device *ptx_dev)
> +fini_streams_for_device (struct ptx_device *ptx_dev)
> {
> free (ptx_dev->async_streams.arr);
> -
> +
> while (ptx_dev->active_streams != NULL)
> {
> - struct PTX_stream *s = ptx_dev->active_streams;
> + struct ptx_stream *s = ptx_dev->active_streams;
> ptx_dev->active_streams = ptx_dev->active_streams->next;
>
> cuStreamDestroy (s->stream);
> map_fini (s);
> free (s);
> }
> -
> +
> map_fini (ptx_dev->null_stream);
> free (ptx_dev->null_stream);
> }
> @@ -463,16 +465,16 @@ fini_streams_for_device (struct PTX_device *ptx_dev)
> associate the stream with the same thread argument. Returns stream to use
> as result. */
>
> -static struct PTX_stream *
> +static struct ptx_stream *
> select_stream_for_async (int async, pthread_t thread, bool create,
> CUstream existing)
> {
> struct nvptx_thread *nvthd = nvptx_thread ();
> /* Local copy of TLS variable. */
> - struct PTX_device *ptx_dev = nvthd->ptx_dev;
> - struct PTX_stream *stream = NULL;
> + struct ptx_device *ptx_dev = nvthd->ptx_dev;
> + struct ptx_stream *stream = NULL;
> int orig_async = async;
> -
> +
> /* The special value acc_async_noval (-1) maps (for now) to an
> implicitly-created stream, which is then handled the same as any other
> numbered async stream. Other options are available, e.g. using the null
> @@ -480,7 +482,7 @@ select_stream_for_async (int async, pthread_t thread, bool create,
> active set. But, stick with this for now. */
> if (async > acc_async_sync)
> async++;
> -
> +
> if (create)
> GOMP_PLUGIN_mutex_lock (&ptx_dev->stream_lock);
>
> @@ -497,19 +499,19 @@ select_stream_for_async (int async, pthread_t thread, bool create,
> else if (async >= 0 && create)
> {
> if (async >= ptx_dev->async_streams.size)
> - {
> + {
> int i, newsize = ptx_dev->async_streams.size * 2;
> -
> +
> if (async >= newsize)
> newsize = async + 1;
> -
> +
> ptx_dev->async_streams.arr
> = GOMP_PLUGIN_realloc (ptx_dev->async_streams.arr,
> - newsize * sizeof (struct PTX_stream *));
> -
> + newsize * sizeof (struct ptx_stream *));
> +
> for (i = ptx_dev->async_streams.size; i < newsize; i++)
> ptx_dev->async_streams.arr[i] = NULL;
> -
> +
> ptx_dev->async_streams.size = newsize;
> }
>
> @@ -519,8 +521,8 @@ select_stream_for_async (int async, pthread_t thread, bool create,
> if (!ptx_dev->async_streams.arr[async] || existing)
> {
> CUresult r;
> - struct PTX_stream *s
> - = GOMP_PLUGIN_malloc (sizeof (struct PTX_stream));
> + struct ptx_stream *s
> + = GOMP_PLUGIN_malloc (sizeof (struct ptx_stream));
>
> if (existing)
> s->stream = existing;
> @@ -528,18 +530,18 @@ select_stream_for_async (int async, pthread_t thread, bool create,
> {
> r = cuStreamCreate (&s->stream, CU_STREAM_DEFAULT);
> if (r != CUDA_SUCCESS)
> - GOMP_PLUGIN_fatal ("cuStreamCreate error: %s", cuErrorMsg (r));
> + GOMP_PLUGIN_fatal ("cuStreamCreate error: %s", cuda_error (r));
> }
> -
> +
> /* If CREATE is true, we're going to be queueing some work on this
> stream. Associate it with the current host thread. */
> s->host_thread = thread;
> s->multithreaded = false;
> -
> +
> s->d = (CUdeviceptr) NULL;
> s->h = NULL;
> map_init (s);
> -
> +
> s->next = ptx_dev->active_streams;
> ptx_dev->active_streams = s;
> ptx_dev->async_streams.arr[async] = s;
> @@ -570,26 +572,20 @@ select_stream_for_async (int async, pthread_t thread, bool create,
> && !pthread_equal (stream->host_thread, thread))
> GOMP_PLUGIN_fatal ("async %d used on wrong thread", orig_async);
>
> -#ifdef DEBUG
> - fprintf (stderr, "libgomp plugin: %s:%s using stream %p (CUDA stream %p) "
> - "for async %d\n", __FILE__, __FUNCTION__, stream,
> - stream ? stream->stream : NULL, orig_async);
> -#endif
> -
> return stream;
> }
>
> -static int PTX_get_num_devices (void);
> +static int nvptx_get_num_devices (void);
>
> /* Initialize the device. */
> static int
> -PTX_init (void)
> +nvptx_init (void)
> {
> CUresult r;
> int rc;
>
> - if (PTX_inited)
> - return PTX_get_num_devices ();
> + if (ptx_inited)
> + return nvptx_get_num_devices ();
>
> rc = verify_device_library ();
> if (rc < 0)
> @@ -597,36 +593,36 @@ PTX_init (void)
>
> r = cuInit (0);
> if (r != CUDA_SUCCESS)
> - GOMP_PLUGIN_fatal ("cuInit error: %s", cuErrorMsg (r));
> + GOMP_PLUGIN_fatal ("cuInit error: %s", cuda_error (r));
>
> - PTX_events = NULL;
> + ptx_events = NULL;
>
> - GOMP_PLUGIN_mutex_init (&PTX_event_lock);
> + GOMP_PLUGIN_mutex_init (&ptx_event_lock);
>
> - PTX_inited = true;
> + ptx_inited = true;
>
> - return PTX_get_num_devices ();
> + return nvptx_get_num_devices ();
> }
>
> static void
> -PTX_fini (void)
> +nvptx_fini (void)
> {
> - PTX_inited = false;
> + ptx_inited = false;
> }
>
> static void *
> -PTX_open_device (int n)
> +nvptx_open_device (int n)
> {
> - struct PTX_device *ptx_dev;
> + struct ptx_device *ptx_dev;
> CUdevice dev;
> CUresult r;
> int async_engines, pi;
>
> r = cuDeviceGet (&dev, n);
> if (r != CUDA_SUCCESS)
> - GOMP_PLUGIN_fatal ("cuDeviceGet error: %s", cuErrorMsg (r));
> + GOMP_PLUGIN_fatal ("cuDeviceGet error: %s", cuda_error (r));
>
> - ptx_dev = GOMP_PLUGIN_malloc (sizeof (struct PTX_device));
> + ptx_dev = GOMP_PLUGIN_malloc (sizeof (struct ptx_device));
>
> ptx_dev->ord = n;
> ptx_dev->dev = dev;
> @@ -634,44 +630,44 @@ PTX_open_device (int n)
>
> r = cuCtxGetCurrent (&ptx_dev->ctx);
> if (r != CUDA_SUCCESS)
> - GOMP_PLUGIN_fatal ("cuCtxGetCurrent error: %s", cuErrorMsg (r));
> + GOMP_PLUGIN_fatal ("cuCtxGetCurrent error: %s", cuda_error (r));
>
> if (!ptx_dev->ctx)
> {
> r = cuCtxCreate (&ptx_dev->ctx, CU_CTX_SCHED_AUTO, dev);
> if (r != CUDA_SUCCESS)
> - GOMP_PLUGIN_fatal ("cuCtxCreate error: %s", cuErrorMsg (r));
> + GOMP_PLUGIN_fatal ("cuCtxCreate error: %s", cuda_error (r));
> }
> else
> ptx_dev->ctx_shared = true;
>
> r = cuDeviceGetAttribute (&pi, CU_DEVICE_ATTRIBUTE_GPU_OVERLAP, dev);
> if (r != CUDA_SUCCESS)
> - GOMP_PLUGIN_fatal ("cuDeviceGetAttribute error: %s", cuErrorMsg (r));
> + GOMP_PLUGIN_fatal ("cuDeviceGetAttribute error: %s", cuda_error (r));
>
> ptx_dev->overlap = pi;
>
> r = cuDeviceGetAttribute (&pi, CU_DEVICE_ATTRIBUTE_CAN_MAP_HOST_MEMORY, dev);
> if (r != CUDA_SUCCESS)
> - GOMP_PLUGIN_fatal ("cuDeviceGetAttribute error: %s", cuErrorMsg (r));
> + GOMP_PLUGIN_fatal ("cuDeviceGetAttribute error: %s", cuda_error (r));
>
> ptx_dev->map = pi;
>
> r = cuDeviceGetAttribute (&pi, CU_DEVICE_ATTRIBUTE_CONCURRENT_KERNELS, dev);
> if (r != CUDA_SUCCESS)
> - GOMP_PLUGIN_fatal ("cuDeviceGetAttribute error: %s", cuErrorMsg (r));
> + GOMP_PLUGIN_fatal ("cuDeviceGetAttribute error: %s", cuda_error (r));
>
> ptx_dev->concur = pi;
>
> r = cuDeviceGetAttribute (&pi, CU_DEVICE_ATTRIBUTE_COMPUTE_MODE, dev);
> if (r != CUDA_SUCCESS)
> - GOMP_PLUGIN_fatal ("cuDeviceGetAttribute error: %s", cuErrorMsg (r));
> + GOMP_PLUGIN_fatal ("cuDeviceGetAttribute error: %s", cuda_error (r));
>
> ptx_dev->mode = pi;
>
> r = cuDeviceGetAttribute (&pi, CU_DEVICE_ATTRIBUTE_INTEGRATED, dev);
> if (r != CUDA_SUCCESS)
> - GOMP_PLUGIN_fatal ("cuDeviceGetAttribute error: %s", cuErrorMsg (r));
> + GOMP_PLUGIN_fatal ("cuDeviceGetAttribute error: %s", cuda_error (r));
>
> ptx_dev->mkern = pi;
>
> @@ -686,21 +682,21 @@ PTX_open_device (int n)
> }
>
> static int
> -PTX_close_device (void *targ_data)
> +nvptx_close_device (void *targ_data)
> {
> CUresult r;
> - struct PTX_device *ptx_dev = targ_data;
> + struct ptx_device *ptx_dev = targ_data;
>
> if (!ptx_dev)
> return 0;
> -
> +
> fini_streams_for_device (ptx_dev);
>
> if (!ptx_dev->ctx_shared)
> {
> r = cuCtxDestroy (ptx_dev->ctx);
> if (r != CUDA_SUCCESS)
> - GOMP_PLUGIN_fatal ("cuCtxDestroy error: %s", cuErrorMsg (r));
> + GOMP_PLUGIN_fatal ("cuCtxDestroy error: %s", cuda_error (r));
> }
>
> free (ptx_dev);
> @@ -709,7 +705,7 @@ PTX_close_device (void *targ_data)
> }
>
> static int
> -PTX_get_num_devices (void)
> +nvptx_get_num_devices (void)
> {
> int n;
> CUresult r;
> @@ -718,16 +714,17 @@ PTX_get_num_devices (void)
> order to enumerate available devices, but CUDA API routines can't be used
> until cuInit has been called. Just call it now (but don't yet do any
> further initialization). */
> - if (!PTX_inited)
> + if (!ptx_inited)
> cuInit (0);
>
> r = cuDeviceGetCount (&n);
> if (r!= CUDA_SUCCESS)
> - GOMP_PLUGIN_fatal ("cuDeviceGetCount error: %s", cuErrorMsg (r));
> + GOMP_PLUGIN_fatal ("cuDeviceGetCount error: %s", cuda_error (r));
>
> return n;
> }
>
> +
> static void
> link_ptx (CUmodule *module, char *ptx_code)
> {
> @@ -743,7 +740,7 @@ link_ptx (CUmodule *module, char *ptx_code)
> void *linkout;
> size_t linkoutsize __attribute__((unused));
>
> - GOMP_PLUGIN_notify ("attempting to load:\n---\n%s\n---\n", ptx_code);
> + GOMP_PLUGIN_debug (0, "attempting to load:\n---\n%s\n---\n", ptx_code);
>
> opts[0] = CU_JIT_WALL_TIME;
> optvals[0] = &elapsed;
> @@ -768,7 +765,7 @@ link_ptx (CUmodule *module, char *ptx_code)
>
> r = cuLinkCreate (7, opts, optvals, &linkstate);
> if (r != CUDA_SUCCESS)
> - GOMP_PLUGIN_fatal ("cuLinkCreate error: %s", cuErrorMsg (r));
> + GOMP_PLUGIN_fatal ("cuLinkCreate error: %s", cuda_error (r));
>
> char *abort_ptx = ABORT_PTX;
> r = cuLinkAddData (linkstate, CU_JIT_INPUT_PTX, abort_ptx,
> @@ -776,7 +773,7 @@ link_ptx (CUmodule *module, char *ptx_code)
> if (r != CUDA_SUCCESS)
> {
> GOMP_PLUGIN_error ("Link error log %s\n", &elog[0]);
> - GOMP_PLUGIN_fatal ("cuLinkAddData (abort) error: %s", cuErrorMsg (r));
> + GOMP_PLUGIN_fatal ("cuLinkAddData (abort) error: %s", cuda_error (r));
> }
>
> char *acc_on_device_ptx = ACC_ON_DEVICE_PTX;
> @@ -786,7 +783,7 @@ link_ptx (CUmodule *module, char *ptx_code)
> {
> GOMP_PLUGIN_error ("Link error log %s\n", &elog[0]);
> GOMP_PLUGIN_fatal ("cuLinkAddData (acc_on_device) error: %s",
> - cuErrorMsg (r));
> + cuda_error (r));
> }
>
> char *goacc_internal_ptx = GOACC_INTERNAL_PTX;
> @@ -796,7 +793,7 @@ link_ptx (CUmodule *module, char *ptx_code)
> {
> GOMP_PLUGIN_error ("Link error log %s\n", &elog[0]);
> GOMP_PLUGIN_fatal ("cuLinkAddData (goacc_internal_ptx) error: %s",
> - cuErrorMsg (r));
> + cuda_error (r));
> }
>
> r = cuLinkAddData (linkstate, CU_JIT_INPUT_PTX, ptx_code,
> @@ -804,33 +801,33 @@ link_ptx (CUmodule *module, char *ptx_code)
> if (r != CUDA_SUCCESS)
> {
> GOMP_PLUGIN_error ("Link error log %s\n", &elog[0]);
> - GOMP_PLUGIN_fatal ("cuLinkAddData (ptx_code) error: %s", cuErrorMsg (r));
> + GOMP_PLUGIN_fatal ("cuLinkAddData (ptx_code) error: %s", cuda_error (r));
> }
>
> r = cuLinkComplete (linkstate, &linkout, &linkoutsize);
> if (r != CUDA_SUCCESS)
> - GOMP_PLUGIN_fatal ("cuLinkComplete error: %s", cuErrorMsg (r));
> + GOMP_PLUGIN_fatal ("cuLinkComplete error: %s", cuda_error (r));
>
> - GOMP_PLUGIN_notify ("Link complete: %fms\n", elapsed);
> - GOMP_PLUGIN_notify ("Link log %s\n", &ilog[0]);
> + GOMP_PLUGIN_debug (0, "Link complete: %fms\n", elapsed);
> + GOMP_PLUGIN_debug (0, "Link log %s\n", &ilog[0]);
>
> r = cuModuleLoadData (module, linkout);
> if (r != CUDA_SUCCESS)
> - GOMP_PLUGIN_fatal ("cuModuleLoadData error: %s", cuErrorMsg (r));
> + GOMP_PLUGIN_fatal ("cuModuleLoadData error: %s", cuda_error (r));
> }
>
> static void
> event_gc (bool memmap_lockable)
> {
> - struct PTX_event *ptx_event = PTX_events;
> + struct ptx_event *ptx_event = ptx_events;
> struct nvptx_thread *nvthd = nvptx_thread ();
>
> - GOMP_PLUGIN_mutex_lock (&PTX_event_lock);
> + GOMP_PLUGIN_mutex_lock (&ptx_event_lock);
>
> while (ptx_event != NULL)
> {
> CUresult r;
> - struct PTX_event *e = ptx_event;
> + struct ptx_event *e = ptx_event;
>
> ptx_event = ptx_event->next;
>
> @@ -849,18 +846,18 @@ event_gc (bool memmap_lockable)
> case PTX_EVT_MEM:
> case PTX_EVT_SYNC:
> break;
> -
> +
> case PTX_EVT_KNL:
> map_pop (e->addr);
> break;
>
> case PTX_EVT_ASYNC_CLEANUP:
> {
> - /* The function GOMP_PLUGIN_async_unmap_vars needs to claim the
> + /* The function gomp_plugin_async_unmap_vars needs to claim the
> memory-map splay tree lock for the current device, so we
> can't call it when one of our callers has already claimed
> the lock. In that case, just delay the GC for this event
> - until later. */
> + until later. */
> if (!memmap_lockable)
> continue;
>
> @@ -872,11 +869,11 @@ event_gc (bool memmap_lockable)
> cuEventDestroy (*te);
> free ((void *)te);
>
> - if (PTX_events == e)
> - PTX_events = PTX_events->next;
> + if (ptx_events == e)
> + ptx_events = ptx_events->next;
> else
> {
> - struct PTX_event *e_ = PTX_events;
> + struct ptx_event *e_ = ptx_events;
> while (e_->next != e)
> e_ = e_->next;
> e_->next = e_->next->next;
> @@ -886,34 +883,34 @@ event_gc (bool memmap_lockable)
> }
> }
>
> - GOMP_PLUGIN_mutex_unlock (&PTX_event_lock);
> + GOMP_PLUGIN_mutex_unlock (&ptx_event_lock);
> }
>
> static void
> -event_add (enum PTX_event_type type, CUevent *e, void *h)
> +event_add (enum ptx_event_type type, CUevent *e, void *h)
> {
> - struct PTX_event *ptx_event;
> + struct ptx_event *ptx_event;
> struct nvptx_thread *nvthd = nvptx_thread ();
>
> assert (type == PTX_EVT_MEM || type == PTX_EVT_KNL || type == PTX_EVT_SYNC
> || type == PTX_EVT_ASYNC_CLEANUP);
>
> - ptx_event = GOMP_PLUGIN_malloc (sizeof (struct PTX_event));
> + ptx_event = GOMP_PLUGIN_malloc (sizeof (struct ptx_event));
> ptx_event->type = type;
> ptx_event->evt = e;
> ptx_event->addr = h;
> ptx_event->ord = nvthd->ptx_dev->ord;
>
> - GOMP_PLUGIN_mutex_lock (&PTX_event_lock);
> + GOMP_PLUGIN_mutex_lock (&ptx_event_lock);
>
> - ptx_event->next = PTX_events;
> - PTX_events = ptx_event;
> + ptx_event->next = ptx_events;
> + ptx_events = ptx_event;
>
> - GOMP_PLUGIN_mutex_unlock (&PTX_event_lock);
> + GOMP_PLUGIN_mutex_unlock (&ptx_event_lock);
> }
>
> void
> -PTX_exec (void (*fn), size_t mapnum, void **hostaddrs, void **devaddrs,
> +nvptx_exec (void (*fn), size_t mapnum, void **hostaddrs, void **devaddrs,
> size_t *sizes, unsigned short *kinds, int num_gangs, int num_workers,
> int vector_length, int async, void *targ_mem_desc)
> {
> @@ -921,14 +918,15 @@ PTX_exec (void (*fn), size_t mapnum, void **hostaddrs, void **devaddrs,
> CUfunction function;
> CUresult r;
> int i;
> - struct PTX_stream *dev_str;
> + struct ptx_stream *dev_str;
> void *kargs[1];
> void *hp, *dp;
> unsigned int nthreads_in_block;
> struct nvptx_thread *nvthd = nvptx_thread ();
> + const char *maybe_abort_msg = "(perhaps abort was called)";
>
> function = targ_fn->fn;
> -
> +
> dev_str = select_stream_for_async (async, pthread_self (), false, NULL);
> assert (dev_str == nvthd->current_stream);
>
> @@ -937,7 +935,7 @@ PTX_exec (void (*fn), size_t mapnum, void **hostaddrs, void **devaddrs,
> the corresponding device pointer. */
> map_push (dev_str, async, mapnum * sizeof (void *), &hp, &dp);
>
> - GOMP_PLUGIN_notify (" %s: prepare mappings\n", __FUNCTION__);
> + GOMP_PLUGIN_debug (0, " %s: prepare mappings\n", __FUNCTION__);
>
> /* Copy the array of arguments to the mapped page. */
> for (i = 0; i < mapnum; i++)
> @@ -947,12 +945,10 @@ PTX_exec (void (*fn), size_t mapnum, void **hostaddrs, void **devaddrs,
> fact have the same value on a unified-memory system). */
> r = cuMemcpy ((CUdeviceptr)dp, (CUdeviceptr)hp, mapnum * sizeof (void *));
> if (r != CUDA_SUCCESS)
> - GOMP_PLUGIN_fatal ("cuMemcpy failed: %s", cuErrorMsg (r));
> + GOMP_PLUGIN_fatal ("cuMemcpy failed: %s", cuda_error (r));
>
> - GOMP_PLUGIN_notify (" %s: kernel %s: launch\n", __FUNCTION__, targ_fn->name);
> + GOMP_PLUGIN_debug (0, " %s: kernel %s: launch\n", __FUNCTION__, targ_fn->name);
>
> - // XXX: possible geometry mappings??
> - //
> // OpenACC CUDA
> //
> // num_gangs blocks
> @@ -978,18 +974,21 @@ PTX_exec (void (*fn), size_t mapnum, void **hostaddrs, void **devaddrs,
>
> kargs[0] = &dp;
> r = cuLaunchKernel (function,
> - num_gangs, 1, 1,
> - nthreads_in_block, 1, 1,
> - 0, dev_str->stream, kargs, 0);
> + num_gangs, 1, 1,
> + nthreads_in_block, 1, 1,
> + 0, dev_str->stream, kargs, 0);
> if (r != CUDA_SUCCESS)
> - GOMP_PLUGIN_fatal ("cuLaunchKernel error: %s", cuErrorMsg (r));
> + GOMP_PLUGIN_fatal ("cuLaunchKernel error: %s", cuda_error (r));
>
> #ifndef DISABLE_ASYNC
> if (async < acc_async_noval)
> {
> r = cuStreamSynchronize (dev_str->stream);
> - if (r != CUDA_SUCCESS)
> - GOMP_PLUGIN_fatal ("cuStreamSynchronize error: %s", cuErrorMsg (r));
> + if (r == CUDA_ERROR_LAUNCH_FAILED)
> + GOMP_PLUGIN_fatal ("cuStreamSynchronize error: %s %s\n", cuda_error (r),
> + maybe_abort_msg);
> + else if (r != CUDA_SUCCESS)
> + GOMP_PLUGIN_fatal ("cuStreamSynchronize error: %s", cuda_error (r));
> }
> else
> {
> @@ -998,25 +997,31 @@ PTX_exec (void (*fn), size_t mapnum, void **hostaddrs, void **devaddrs,
> e = (CUevent *)GOMP_PLUGIN_malloc (sizeof (CUevent));
>
> r = cuEventCreate (e, CU_EVENT_DISABLE_TIMING);
> - if (r != CUDA_SUCCESS)
> - GOMP_PLUGIN_fatal ("cuEventCreate error: %s", cuErrorMsg (r));
> + if (r == CUDA_ERROR_LAUNCH_FAILED)
> + GOMP_PLUGIN_fatal ("cuEventCreate error: %s %s\n", cuda_error (r),
> + maybe_abort_msg);
> + else if (r != CUDA_SUCCESS)
> + GOMP_PLUGIN_fatal ("cuEventCreate error: %s", cuda_error (r));
>
> event_gc (true);
>
> r = cuEventRecord (*e, dev_str->stream);
> if (r != CUDA_SUCCESS)
> - GOMP_PLUGIN_fatal ("cuEventRecord error: %s", cuErrorMsg (r));
> + GOMP_PLUGIN_fatal ("cuEventRecord error: %s", cuda_error (r));
>
> event_add (PTX_EVT_KNL, e, (void *)dev_str);
> }
> #else
> r = cuCtxSynchronize ();
> - if (r != CUDA_SUCCESS)
> - GOMP_PLUGIN_fatal ("cuCtxSynchronize error: %s", cuErrorMsg (r));
> + if (r == CUDA_ERROR_LAUNCH_FAILED)
> + GOMP_PLUGIN_fatal ("cuCtxSynchronize error: %s %s\n", cuda_error (r),
> + maybe_abort_msg);
> + else if (r != CUDA_SUCCESS)
> + GOMP_PLUGIN_fatal ("cuCtxSynchronize error: %s", cuda_error (r));
> #endif
>
> - GOMP_PLUGIN_notify (" %s: kernel %s: finished\n", __FUNCTION__,
> - targ_fn->name);
> + GOMP_PLUGIN_debug (0, " %s: kernel %s: finished\n", __FUNCTION__,
> + targ_fn->name);
>
> #ifndef DISABLE_ASYNC
> if (async < acc_async_noval)
> @@ -1027,7 +1032,7 @@ PTX_exec (void (*fn), size_t mapnum, void **hostaddrs, void **devaddrs,
> void * openacc_get_current_cuda_context (void);
>
> static void *
> -PTX_alloc (size_t s)
> +nvptx_alloc (size_t s)
> {
> CUdeviceptr d;
> CUresult r;
> @@ -1036,12 +1041,12 @@ PTX_alloc (size_t s)
> if (r == CUDA_ERROR_OUT_OF_MEMORY)
> return 0;
> if (r != CUDA_SUCCESS)
> - GOMP_PLUGIN_fatal ("cuMemAlloc error: %s", cuErrorMsg (r));
> + GOMP_PLUGIN_fatal ("cuMemAlloc error: %s", cuda_error (r));
> return (void *)d;
> }
>
> static void
> -PTX_free (void *p)
> +nvptx_free (void *p)
> {
> CUresult r;
> CUdeviceptr pb;
> @@ -1049,18 +1054,18 @@ PTX_free (void *p)
>
> r = cuMemGetAddressRange (&pb, &ps, (CUdeviceptr)p);
> if (r != CUDA_SUCCESS)
> - GOMP_PLUGIN_fatal ("cuMemGetAddressRange error: %s", cuErrorMsg (r));
> + GOMP_PLUGIN_fatal ("cuMemGetAddressRange error: %s", cuda_error (r));
>
> if ((CUdeviceptr)p != pb)
> GOMP_PLUGIN_fatal ("invalid device address");
>
> r = cuMemFree ((CUdeviceptr)p);
> if (r != CUDA_SUCCESS)
> - GOMP_PLUGIN_fatal ("cuMemFree error: %s", cuErrorMsg (r));
> + GOMP_PLUGIN_fatal ("cuMemFree error: %s", cuda_error (r));
> }
>
> static void *
> -PTX_host2dev (void *d, const void *h, size_t s)
> +nvptx_host2dev (void *d, const void *h, size_t s)
> {
> CUresult r;
> CUdeviceptr pb;
> @@ -1075,7 +1080,7 @@ PTX_host2dev (void *d, const void *h, size_t s)
>
> r = cuMemGetAddressRange (&pb, &ps, (CUdeviceptr)d);
> if (r != CUDA_SUCCESS)
> - GOMP_PLUGIN_fatal ("cuMemGetAddressRange error: %s", cuErrorMsg (r));
> + GOMP_PLUGIN_fatal ("cuMemGetAddressRange error: %s", cuda_error (r));
>
> if (!pb)
> GOMP_PLUGIN_fatal ("invalid device address");
> @@ -1098,18 +1103,18 @@ PTX_host2dev (void *d, const void *h, size_t s)
>
> r = cuEventCreate (e, CU_EVENT_DISABLE_TIMING);
> if (r != CUDA_SUCCESS)
> - GOMP_PLUGIN_fatal ("cuEventCreate error: %s", cuErrorMsg (r));
> + GOMP_PLUGIN_fatal ("cuEventCreate error: %s", cuda_error (r));
>
> event_gc (false);
>
> r = cuMemcpyHtoDAsync ((CUdeviceptr)d, h, s,
> nvthd->current_stream->stream);
> if (r != CUDA_SUCCESS)
> - GOMP_PLUGIN_fatal ("cuMemcpyHtoDAsync error: %s", cuErrorMsg (r));
> + GOMP_PLUGIN_fatal ("cuMemcpyHtoDAsync error: %s", cuda_error (r));
>
> r = cuEventRecord (*e, nvthd->current_stream->stream);
> if (r != CUDA_SUCCESS)
> - GOMP_PLUGIN_fatal ("cuEventRecord error: %s", cuErrorMsg (r));
> + GOMP_PLUGIN_fatal ("cuEventRecord error: %s", cuda_error (r));
>
> event_add (PTX_EVT_MEM, e, (void *)h);
> }
> @@ -1118,14 +1123,14 @@ PTX_host2dev (void *d, const void *h, size_t s)
> {
> r = cuMemcpyHtoD ((CUdeviceptr)d, h, s);
> if (r != CUDA_SUCCESS)
> - GOMP_PLUGIN_fatal ("cuMemcpyHtoD error: %s", cuErrorMsg (r));
> + GOMP_PLUGIN_fatal ("cuMemcpyHtoD error: %s", cuda_error (r));
> }
>
> return 0;
> }
>
> static void *
> -PTX_dev2host (void *h, const void *d, size_t s)
> +nvptx_dev2host (void *h, const void *d, size_t s)
> {
> CUresult r;
> CUdeviceptr pb;
> @@ -1140,7 +1145,7 @@ PTX_dev2host (void *h, const void *d, size_t s)
>
> r = cuMemGetAddressRange (&pb, &ps, (CUdeviceptr)d);
> if (r != CUDA_SUCCESS)
> - GOMP_PLUGIN_fatal ("cuMemGetAddressRange error: %s", cuErrorMsg (r));
> + GOMP_PLUGIN_fatal ("cuMemGetAddressRange error: %s", cuda_error (r));
>
> if (!pb)
> GOMP_PLUGIN_fatal ("invalid device address");
> @@ -1163,18 +1168,18 @@ PTX_dev2host (void *h, const void *d, size_t s)
>
> r = cuEventCreate (e, CU_EVENT_DISABLE_TIMING);
> if (r != CUDA_SUCCESS)
> - GOMP_PLUGIN_fatal ("cuEventCreate error: %s\n", cuErrorMsg (r));
> + GOMP_PLUGIN_fatal ("cuEventCreate error: %s\n", cuda_error (r));
>
> event_gc (false);
>
> r = cuMemcpyDtoHAsync (h, (CUdeviceptr)d, s,
> nvthd->current_stream->stream);
> if (r != CUDA_SUCCESS)
> - GOMP_PLUGIN_fatal ("cuMemcpyDtoHAsync error: %s", cuErrorMsg (r));
> + GOMP_PLUGIN_fatal ("cuMemcpyDtoHAsync error: %s", cuda_error (r));
>
> r = cuEventRecord (*e, nvthd->current_stream->stream);
> if (r != CUDA_SUCCESS)
> - GOMP_PLUGIN_fatal ("cuEventRecord error: %s", cuErrorMsg (r));
> + GOMP_PLUGIN_fatal ("cuEventRecord error: %s", cuda_error (r));
>
> event_add (PTX_EVT_MEM, e, (void *)h);
> }
> @@ -1183,14 +1188,14 @@ PTX_dev2host (void *h, const void *d, size_t s)
> {
> r = cuMemcpyDtoH (h, (CUdeviceptr)d, s);
> if (r != CUDA_SUCCESS)
> - GOMP_PLUGIN_fatal ("cuMemcpyDtoH error: %s", cuErrorMsg (r));
> + GOMP_PLUGIN_fatal ("cuMemcpyDtoH error: %s", cuda_error (r));
> }
>
> return 0;
> }
>
> static void
> -PTX_set_async (int async)
> +nvptx_set_async (int async)
> {
> struct nvptx_thread *nvthd = nvptx_thread ();
> nvthd->current_stream
> @@ -1198,11 +1203,11 @@ PTX_set_async (int async)
> }
>
> static int
> -PTX_async_test (int async)
> +nvptx_async_test (int async)
> {
> CUresult r;
> - struct PTX_stream *s;
> -
> + struct ptx_stream *s;
> +
> s = select_stream_for_async (async, pthread_self (), false, NULL);
>
> if (!s)
> @@ -1222,15 +1227,15 @@ PTX_async_test (int async)
> else if (r == CUDA_ERROR_NOT_READY)
> return 0;
>
> - GOMP_PLUGIN_fatal ("cuStreamQuery error: %s", cuErrorMsg (r));
> + GOMP_PLUGIN_fatal ("cuStreamQuery error: %s", cuda_error (r));
>
> return 0;
> }
>
> static int
> -PTX_async_test_all (void)
> +nvptx_async_test_all (void)
> {
> - struct PTX_stream *s;
> + struct ptx_stream *s;
> pthread_t self = pthread_self ();
> struct nvptx_thread *nvthd = nvptx_thread ();
>
> @@ -1254,11 +1259,11 @@ PTX_async_test_all (void)
> }
>
> static void
> -PTX_wait (int async)
> +nvptx_wait (int async)
> {
> CUresult r;
> - struct PTX_stream *s;
> -
> + struct ptx_stream *s;
> +
> s = select_stream_for_async (async, pthread_self (), false, NULL);
>
> if (!s)
> @@ -1266,17 +1271,17 @@ PTX_wait (int async)
>
> r = cuStreamSynchronize (s->stream);
> if (r != CUDA_SUCCESS)
> - GOMP_PLUGIN_fatal ("cuStreamSynchronize error: %s", cuErrorMsg (r));
> -
> + GOMP_PLUGIN_fatal ("cuStreamSynchronize error: %s", cuda_error (r));
> +
> event_gc (true);
> }
>
> static void
> -PTX_wait_async (int async1, int async2)
> +nvptx_wait_async (int async1, int async2)
> {
> CUresult r;
> CUevent *e;
> - struct PTX_stream *s1, *s2;
> + struct ptx_stream *s1, *s2;
> pthread_t self = pthread_self ();
>
> /* The stream that is waiting (rather than being waited for) doesn't
> @@ -1294,26 +1299,26 @@ PTX_wait_async (int async1, int async2)
>
> r = cuEventCreate (e, CU_EVENT_DISABLE_TIMING);
> if (r != CUDA_SUCCESS)
> - GOMP_PLUGIN_fatal ("cuEventCreate error: %s", cuErrorMsg (r));
> + GOMP_PLUGIN_fatal ("cuEventCreate error: %s", cuda_error (r));
>
> event_gc (true);
>
> r = cuEventRecord (*e, s1->stream);
> if (r != CUDA_SUCCESS)
> - GOMP_PLUGIN_fatal ("cuEventRecord error: %s", cuErrorMsg (r));
> + GOMP_PLUGIN_fatal ("cuEventRecord error: %s", cuda_error (r));
>
> event_add (PTX_EVT_SYNC, e, NULL);
>
> r = cuStreamWaitEvent (s2->stream, *e, 0);
> if (r != CUDA_SUCCESS)
> - GOMP_PLUGIN_fatal ("cuStreamWaitEvent error: %s", cuErrorMsg (r));
> + GOMP_PLUGIN_fatal ("cuStreamWaitEvent error: %s", cuda_error (r));
> }
>
> static void
> -PTX_wait_all (void)
> +nvptx_wait_all (void)
> {
> CUresult r;
> - struct PTX_stream *s;
> + struct ptx_stream *s;
> pthread_t self = pthread_self ();
> struct nvptx_thread *nvthd = nvptx_thread ();
>
> @@ -1324,16 +1329,16 @@ PTX_wait_all (void)
> for (s = nvthd->ptx_dev->active_streams; s != NULL; s = s->next)
> {
> if (s->multithreaded || pthread_equal (s->host_thread, self))
> - {
> + {
> r = cuStreamQuery (s->stream);
> if (r == CUDA_SUCCESS)
> continue;
> else if (r != CUDA_ERROR_NOT_READY)
> - GOMP_PLUGIN_fatal ("cuStreamQuery error: %s", cuErrorMsg (r));
> + GOMP_PLUGIN_fatal ("cuStreamQuery error: %s", cuda_error (r));
>
> r = cuStreamSynchronize (s->stream);
> if (r != CUDA_SUCCESS)
> - GOMP_PLUGIN_fatal ("cuStreamSynchronize error: %s", cuErrorMsg (r));
> + GOMP_PLUGIN_fatal ("cuStreamSynchronize error: %s", cuda_error (r));
> }
> }
>
> @@ -1343,19 +1348,19 @@ PTX_wait_all (void)
> }
>
> static void
> -PTX_wait_all_async (int async)
> +nvptx_wait_all_async (int async)
> {
> CUresult r;
> - struct PTX_stream *waiting_stream, *other_stream;
> + struct ptx_stream *waiting_stream, *other_stream;
> CUevent *e;
> struct nvptx_thread *nvthd = nvptx_thread ();
> pthread_t self = pthread_self ();
> -
> +
> /* The stream doing the waiting. This could be the first mention of the
> stream, so create it if necessary. */
> waiting_stream
> = select_stream_for_async (async, pthread_self (), true, NULL);
> -
> +
> /* Launches on the null stream already block on other streams in the
> context. */
> if (!waiting_stream || waiting_stream == nvthd->ptx_dev->null_stream)
> @@ -1377,25 +1382,25 @@ PTX_wait_all_async (int async)
>
> r = cuEventCreate (e, CU_EVENT_DISABLE_TIMING);
> if (r != CUDA_SUCCESS)
> - GOMP_PLUGIN_fatal ("cuEventCreate error: %s", cuErrorMsg (r));
> + GOMP_PLUGIN_fatal ("cuEventCreate error: %s", cuda_error (r));
>
> /* Record an event on the waited-for stream. */
> r = cuEventRecord (*e, other_stream->stream);
> if (r != CUDA_SUCCESS)
> - GOMP_PLUGIN_fatal ("cuEventRecord error: %s", cuErrorMsg (r));
> + GOMP_PLUGIN_fatal ("cuEventRecord error: %s", cuda_error (r));
>
> event_add (PTX_EVT_SYNC, e, NULL);
>
> r = cuStreamWaitEvent (waiting_stream->stream, *e, 0);
> if (r != CUDA_SUCCESS)
> - GOMP_PLUGIN_fatal ("cuStreamWaitEvent error: %s", cuErrorMsg (r));
> + GOMP_PLUGIN_fatal ("cuStreamWaitEvent error: %s", cuda_error (r));
> }
>
> GOMP_PLUGIN_mutex_unlock (&nvthd->ptx_dev->stream_lock);
> }
>
> static void *
> -PTX_get_current_cuda_device (void)
> +nvptx_get_current_cuda_device (void)
> {
> struct nvptx_thread *nvthd = nvptx_thread ();
>
> @@ -1406,7 +1411,7 @@ PTX_get_current_cuda_device (void)
> }
>
> static void *
> -PTX_get_current_cuda_context (void)
> +nvptx_get_current_cuda_context (void)
> {
> struct nvptx_thread *nvthd = nvptx_thread ();
>
> @@ -1417,9 +1422,9 @@ PTX_get_current_cuda_context (void)
> }
>
> static void *
> -PTX_get_cuda_stream (int async)
> +nvptx_get_cuda_stream (int async)
> {
> - struct PTX_stream *s;
> + struct ptx_stream *s;
> struct nvptx_thread *nvthd = nvptx_thread ();
>
> if (!nvthd || !nvthd->ptx_dev)
> @@ -1431,9 +1436,9 @@ PTX_get_cuda_stream (int async)
> }
>
> static int
> -PTX_set_cuda_stream (int async, void *stream)
> +nvptx_set_cuda_stream (int async, void *stream)
> {
> - struct PTX_stream *oldstream;
> + struct ptx_stream *oldstream;
> pthread_t self = pthread_self ();
> struct nvptx_thread *nvthd = nvptx_thread ();
>
> @@ -1451,14 +1456,14 @@ PTX_set_cuda_stream (int async, void *stream)
> returned from acc_get_cuda_stream above... */
>
> oldstream = select_stream_for_async (async, self, false, NULL);
> -
> +
> if (oldstream)
> {
> if (nvthd->ptx_dev->active_streams == oldstream)
> nvthd->ptx_dev->active_streams = nvthd->ptx_dev->active_streams->next;
> else
> {
> - struct PTX_stream *s = nvthd->ptx_dev->active_streams;
> + struct ptx_stream *s = nvthd->ptx_dev->active_streams;
> while (s->next != oldstream)
> s = s->next;
> s->next = s->next->next;
> @@ -1482,10 +1487,6 @@ PTX_set_cuda_stream (int async, void *stream)
> int
> GOMP_OFFLOAD_get_type (void)
> {
> -#ifdef DEBUG
> - fprintf (stderr, "libgomp plugin: %s:%s\n", __FILE__, __FUNCTION__);
> -#endif
> -
> return OFFLOAD_TARGET_TYPE_NVIDIA_PTX;
> }
>
> @@ -1504,11 +1505,7 @@ GOMP_OFFLOAD_get_name (void)
> int
> GOMP_OFFLOAD_get_num_devices (void)
> {
> -#ifdef DEBUG
> - fprintf (stderr, "libgomp plugin: %s:%s\n", __FILE__, __FUNCTION__);
> -#endif
> -
> - return PTX_get_num_devices ();
> + return nvptx_get_num_devices ();
> }
>
> static void **kernel_target_data;
> @@ -1517,11 +1514,6 @@ static void **kernel_host_table;
> void
> GOMP_OFFLOAD_register_image (void *host_table, void *target_data)
> {
> -#ifdef DEBUG
> - fprintf (stderr, "libgomp plugin: %s:%s (%p, %p)\n", __FILE__, __FUNCTION__,
> - host_table, target_data);
> -#endif
> -
> kernel_target_data = target_data;
> kernel_host_table = host_table;
> }
> @@ -1529,21 +1521,13 @@ GOMP_OFFLOAD_register_image (void *host_table, void *target_data)
> void
> GOMP_OFFLOAD_init_device (int n __attribute__((unused)))
> {
> -#ifdef DEBUG
> - fprintf (stderr, "libgomp plugin: %s:%s\n", __FILE__, __FUNCTION__);
> -#endif
> -
> - (void) PTX_init ();
> + (void) nvptx_init ();
> }
>
> void
> GOMP_OFFLOAD_fini_device (int n __attribute__((unused)))
> {
> -#ifdef DEBUG
> - fprintf (stderr, "libgomp plugin: %s:%s\n", __FILE__, __FUNCTION__);
> -#endif
> -
> - PTX_fini ();
> + nvptx_fini ();
> }
>
> int
> @@ -1557,12 +1541,7 @@ GOMP_OFFLOAD_get_table (int n __attribute__((unused)),
> CUresult r;
> struct targ_fn_descriptor *targ_fns;
>
> -#ifdef DEBUG
> - fprintf (stderr, "libgomp plugin: %s:%s (%p)\n", __FILE__, __FUNCTION__,
> - tablep);
> -#endif
> -
> - if (PTX_init () <= 0)
> + if (nvptx_init () <= 0)
> return 0;
>
> /* This isn't an error, because an image may legitimately have no offloaded
> @@ -1596,11 +1575,11 @@ GOMP_OFFLOAD_get_table (int n __attribute__((unused)),
>
> r = cuModuleGetFunction (&function, module, fn_names[i]);
> if (r != CUDA_SUCCESS)
> - GOMP_PLUGIN_fatal ("cuModuleGetFunction error: %s", cuErrorMsg (r));
> + GOMP_PLUGIN_fatal ("cuModuleGetFunction error: %s", cuda_error (r));
>
> targ_fns[i].fn = function;
> targ_fns[i].name = (const char *) fn_names[i];
> -
> +
> (*tablep)[i].host_start = (uintptr_t) fn_table[i];
> (*tablep)[i].host_end = (*tablep)[i].host_start + 1;
> (*tablep)[i].tgt_start = (uintptr_t) &targ_fns[i];
> @@ -1613,84 +1592,52 @@ GOMP_OFFLOAD_get_table (int n __attribute__((unused)),
> void *
> GOMP_OFFLOAD_alloc (int n __attribute__((unused)), size_t size)
> {
> -#ifdef DEBUG
> - fprintf (stderr, "libgomp plugin: %s:%s (%zu)\n", __FILE__, __FUNCTION__,
> - size);
> -#endif
> -
> - return PTX_alloc (size);
> + return nvptx_alloc (size);
> }
>
> void
> GOMP_OFFLOAD_free (int n __attribute__((unused)), void *ptr)
> {
> -#ifdef DEBUG
> - fprintf (stderr, "libgomp plugin: %s:%s (%p)\n", __FILE__, __FUNCTION__, ptr);
> -#endif
> -
> - PTX_free (ptr);
> + nvptx_free (ptr);
> }
>
> void *
> GOMP_OFFLOAD_dev2host (int ord __attribute__((unused)), void *dst,
> const void *src, size_t n)
> {
> -#ifdef DEBUG
> - fprintf (stderr, "libgomp plugin: %s:%s (%p, %p, %zu)\n", __FILE__,
> - __FUNCTION__, dst,
> - src, n);
> -#endif
> -
> - return PTX_dev2host (dst, src, n);
> + return nvptx_dev2host (dst, src, n);
> }
>
> void *
> GOMP_OFFLOAD_host2dev (int ord __attribute__((unused)), void *dst,
> const void *src, size_t n)
> {
> -#ifdef DEBUG
> - fprintf (stderr, "libgomp plugin: %s:%s (%p, %p, %zu)\n", __FILE__,
> - __FUNCTION__, dst, src, n);
> -#endif
> -
> - return PTX_host2dev (dst, src, n);
> + return nvptx_host2dev (dst, src, n);
> }
>
> void (*device_run) (int n, void *fn_ptr, void *vars) = NULL;
>
> void
> GOMP_OFFLOAD_openacc_parallel (void (*fn) (void *), size_t mapnum,
> - void **hostaddrs, void **devaddrs, size_t *sizes,
> - unsigned short *kinds, int num_gangs,
> - int num_workers, int vector_length, int async,
> - void *targ_mem_desc)
> + void **hostaddrs, void **devaddrs, size_t *sizes,
> + unsigned short *kinds, int num_gangs,
> + int num_workers, int vector_length, int async,
> + void *targ_mem_desc)
> {
> -#ifdef DEBUG
> - fprintf (stderr, "libgomp plugin: %s:%s (%p, %zu, %p, %p, %p, %d, %d, %d, "
> - "%d, %p)\n", __FILE__, __FUNCTION__, fn, mapnum, hostaddrs, sizes,
> - kinds, num_gangs, num_workers, vector_length, async, targ_mem_desc);
> -#endif
> -
> - PTX_exec (fn, mapnum, hostaddrs, devaddrs, sizes, kinds, num_gangs,
> + nvptx_exec (fn, mapnum, hostaddrs, devaddrs, sizes, kinds, num_gangs,
> num_workers, vector_length, async, targ_mem_desc);
> }
>
> void *
> GOMP_OFFLOAD_openacc_open_device (int n)
> {
> -#ifdef DEBUG
> - fprintf (stderr, "libgomp plugin: %s:%s (%d)\n", __FILE__, __FUNCTION__, n);
> -#endif
> - return PTX_open_device (n);
> + return nvptx_open_device (n);
> }
>
> int
> GOMP_OFFLOAD_openacc_close_device (void *h)
> {
> -#ifdef DEBUG
> - fprintf (stderr, "libgomp plugin: %s:%s (%p)\n", __FILE__, __FUNCTION__, h);
> -#endif
> - return PTX_close_device (h);
> + return nvptx_close_device (h);
> }
>
> void
> @@ -1701,7 +1648,7 @@ GOMP_OFFLOAD_openacc_set_device_num (int n)
> assert (n >= 0);
>
> if (!nvthd->ptx_dev || nvthd->ptx_dev->ord != n)
> - (void) PTX_open_device (n);
> + (void) nvptx_open_device (n);
> }
>
> /* This can be called before the device is "opened" for the current thread, in
> @@ -1727,20 +1674,15 @@ GOMP_OFFLOAD_openacc_register_async_cleanup (void *targ_mem_desc)
> CUresult r;
> struct nvptx_thread *nvthd = nvptx_thread ();
>
> -#ifdef DEBUG
> - fprintf (stderr, "libgomp plugin: %s:%s (%p)\n", __FILE__, __FUNCTION__,
> - targ_mem_desc);
> -#endif
> -
> e = (CUevent *) GOMP_PLUGIN_malloc (sizeof (CUevent));
>
> r = cuEventCreate (e, CU_EVENT_DISABLE_TIMING);
> if (r != CUDA_SUCCESS)
> - GOMP_PLUGIN_fatal ("cuEventCreate error: %s", cuErrorMsg (r));
> + GOMP_PLUGIN_fatal ("cuEventCreate error: %s", cuda_error (r));
>
> r = cuEventRecord (*e, nvthd->current_stream->stream);
> if (r != CUDA_SUCCESS)
> - GOMP_PLUGIN_fatal ("cuEventRecord error: %s", cuErrorMsg (r));
> + GOMP_PLUGIN_fatal ("cuEventRecord error: %s", cuda_error (r));
>
> event_add (PTX_EVT_ASYNC_CLEANUP, e, targ_mem_desc);
> }
> @@ -1748,75 +1690,49 @@ GOMP_OFFLOAD_openacc_register_async_cleanup (void *targ_mem_desc)
> int
> GOMP_OFFLOAD_openacc_async_test (int async)
> {
> -#ifdef DEBUG
> - fprintf (stderr, "libgomp plugin: %s:%s (%d)\n", __FILE__, __FUNCTION__,
> - async);
> -#endif
> - return PTX_async_test (async);
> + return nvptx_async_test (async);
> }
>
> int
> GOMP_OFFLOAD_openacc_async_test_all (void)
> {
> -#ifdef DEBUG
> - fprintf (stderr, "libgomp plugin: %s:%s\n", __FILE__, __FUNCTION__);
> -#endif
> - return PTX_async_test_all ();
> + return nvptx_async_test_all ();
> }
>
> void
> GOMP_OFFLOAD_openacc_async_wait (int async)
> {
> -#ifdef DEBUG
> - fprintf (stderr, "libgomp plugin: %s:%s (%d)\n", __FILE__, __FUNCTION__,
> - async);
> -#endif
> - PTX_wait (async);
> + nvptx_wait (async);
> }
>
> void
> GOMP_OFFLOAD_openacc_async_wait_async (int async1, int async2)
> {
> -#ifdef DEBUG
> - fprintf (stderr, "libgomp plugin: %s:%s (%d, %d)\n", __FILE__, __FUNCTION__,
> - async1, async2);
> -#endif
> - PTX_wait_async (async1, async2);
> + nvptx_wait_async (async1, async2);
> }
>
> void
> GOMP_OFFLOAD_openacc_async_wait_all (void)
> {
> -#ifdef DEBUG
> - fprintf (stderr, "libgomp plugin: %s:%s\n", __FILE__, __FUNCTION__);
> -#endif
> - PTX_wait_all ();
> + nvptx_wait_all ();
> }
>
> void
> GOMP_OFFLOAD_openacc_async_wait_all_async (int async)
> {
> -#ifdef DEBUG
> - fprintf (stderr, "libgomp plugin: %s:%s (%d)\n", __FILE__, __FUNCTION__,
> - async);
> -#endif
> - PTX_wait_all_async (async);
> + nvptx_wait_all_async (async);
> }
>
> void
> GOMP_OFFLOAD_openacc_async_set_async (int async)
> {
> -#ifdef DEBUG
> - fprintf (stderr, "libgomp plugin: %s:%s (%d)\n", __FILE__, __FUNCTION__,
> - async);
> -#endif
> - PTX_set_async (async);
> + nvptx_set_async (async);
> }
>
> void *
> GOMP_OFFLOAD_openacc_create_thread_data (void *targ_data)
> {
> - struct PTX_device *ptx_dev = (struct PTX_device *) targ_data;
> + struct ptx_device *ptx_dev = (struct ptx_device *) targ_data;
> struct nvptx_thread *nvthd
> = GOMP_PLUGIN_malloc (sizeof (struct nvptx_thread));
> CUresult r;
> @@ -1824,7 +1740,7 @@ GOMP_OFFLOAD_openacc_create_thread_data (void *targ_data)
>
> r = cuCtxGetCurrent (&thd_ctx);
> if (r != CUDA_SUCCESS)
> - GOMP_PLUGIN_fatal ("cuCtxGetCurrent error: %s", cuErrorMsg (r));
> + GOMP_PLUGIN_fatal ("cuCtxGetCurrent error: %s", cuda_error (r));
>
> assert (ptx_dev->ctx);
>
> @@ -1832,7 +1748,7 @@ GOMP_OFFLOAD_openacc_create_thread_data (void *targ_data)
> {
> r = cuCtxPushCurrent (ptx_dev->ctx);
> if (r != CUDA_SUCCESS)
> - GOMP_PLUGIN_fatal ("cuCtxPushCurrent error: %s", cuErrorMsg (r));
> + GOMP_PLUGIN_fatal ("cuCtxPushCurrent error: %s", cuda_error (r));
> }
>
> nvthd->current_stream = ptx_dev->null_stream;
> @@ -1850,41 +1766,27 @@ GOMP_OFFLOAD_openacc_destroy_thread_data (void *data)
> void *
> GOMP_OFFLOAD_openacc_get_current_cuda_device (void)
> {
> -#ifdef DEBUG
> - fprintf (stderr, "libgomp plugin: %s:%s\n", __FILE__, __FUNCTION__);
> -#endif
> - return PTX_get_current_cuda_device ();
> + return nvptx_get_current_cuda_device ();
> }
>
> void *
> GOMP_OFFLOAD_openacc_get_current_cuda_context (void)
> {
> -#ifdef DEBUG
> - fprintf (stderr, "libgomp plugin: %s:%s\n", __FILE__, __FUNCTION__);
> -#endif
> - return PTX_get_current_cuda_context ();
> + return nvptx_get_current_cuda_context ();
> }
>
> -/* NOTE: This returns a CUstream, not a PTX_stream pointer. */
> +/* NOTE: This returns a CUstream, not a ptx_stream pointer. */
>
> void *
> GOMP_OFFLOAD_openacc_get_cuda_stream (int async)
> {
> -#ifdef DEBUG
> - fprintf (stderr, "libgomp plugin: %s:%s (%d)\n", __FILE__, __FUNCTION__,
> - async);
> -#endif
> - return PTX_get_cuda_stream (async);
> + return nvptx_get_cuda_stream (async);
> }
>
> -/* NOTE: This takes a CUstream, not a PTX_stream pointer. */
> +/* NOTE: This takes a CUstream, not a ptx_stream pointer. */
>
> int
> GOMP_OFFLOAD_openacc_set_cuda_stream (int async, void *stream)
> {
> -#ifdef DEBUG
> - fprintf (stderr, "libgomp plugin: %s:%s (%d, %p)\n", __FILE__, __FUNCTION__,
> - async, stream);
> -#endif
> - return PTX_set_cuda_stream (async, stream);
> + return nvptx_set_cuda_stream (async, stream);
> }
> diff --git libgomp/sections.c libgomp/sections.c
> index fb746c7..654daa9 100644
> --- libgomp/sections.c
> +++ libgomp/sections.c
> @@ -1,7 +1,8 @@
> /* Copyright (C) 2005-2014 Free Software Foundation, Inc.
> Contributed by Richard Henderson <rth@redhat.com>.
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> diff --git libgomp/single.c libgomp/single.c
> index afa94fd..88c9fbe 100644
> --- libgomp/single.c
> +++ libgomp/single.c
> @@ -1,7 +1,8 @@
> /* Copyright (C) 2005-2014 Free Software Foundation, Inc.
> Contributed by Richard Henderson <rth@redhat.com>.
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> diff --git libgomp/splay-tree.c libgomp/splay-tree.c
> index 14b03ac..25532ae 100644
> --- libgomp/splay-tree.c
> +++ libgomp/splay-tree.c
> @@ -3,7 +3,8 @@
> Free Software Foundation, Inc.
> Contributed by Mark Mitchell (mark@markmitchell.com).
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> diff --git libgomp/splay-tree.h libgomp/splay-tree.h
> index f29d437..2afbbc2 100644
> --- libgomp/splay-tree.h
> +++ libgomp/splay-tree.h
> @@ -3,7 +3,8 @@
> Free Software Foundation, Inc.
> Contributed by Mark Mitchell (mark@markmitchell.com).
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> diff --git libgomp/target.c libgomp/target.c
> index 93fd4f4..c92b43a 100644
> --- libgomp/target.c
> +++ libgomp/target.c
> @@ -1,7 +1,8 @@
> /* Copyright (C) 2013-2014 Free Software Foundation, Inc.
> Contributed by Jakub Jelinek <jakub@redhat.com>.
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> @@ -436,11 +437,11 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
>
> for (j = i + 1; j < mapnum; j++)
> if (!GOMP_MAP_POINTER_P (get_kind (is_openacc, kinds, j)
> - & typemask))
> + & typemask))
> break;
> else if ((uintptr_t) hostaddrs[j] < k->host_start
> - || ((uintptr_t) hostaddrs[j] + sizeof (void *)
> - > k->host_end))
> + || ((uintptr_t) hostaddrs[j] + sizeof (void *)
> + > k->host_end))
> break;
> else
> {
> @@ -505,34 +506,30 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
> sizeof (void *));
> i++;
> }
> - break;
> - }
> - case GOMP_MAP_FORCE_PRESENT:
> - {
> - /* We already looked up the memory region above and it
> - was missing. */
> - size_t size = k->host_end - k->host_start;
> - gomp_fatal ("present clause: !acc_is_present (%p, "
> - "%zd (0x%zx))", (void *) k->host_start,
> - size, size);
> - }
> - break;
> - case GOMP_MAP_FORCE_DEVICEPTR:
> - assert (k->host_end - k->host_start == sizeof (void *));
> -
> - devicep->host2dev_func (devicep->target_id,
> - (void *) (tgt->tgt_start
> - + k->tgt_offset),
> - (void *) k->host_start,
> - sizeof (void *));
> - break;
> - case GOMP_MAP_FORCE_PRIVATE:
> - abort ();
> - case GOMP_MAP_FORCE_FIRSTPRIVATE:
> - abort ();
> - default:
> - gomp_fatal ("%s: unhandled kind 0x%.2x", __FUNCTION__,
> - kind);
> + }
> + break;
> + case GOMP_MAP_FORCE_PRESENT:
> + {
> + /* We already looked up the memory region above and it
> + was missing. */
> + size_t size = k->host_end - k->host_start;
> + gomp_fatal ("present clause: !acc_is_present (%p, "
> + "%zd (0x%zx))", (void *) k->host_start,
> + size, size);
> + }
> + break;
> + case GOMP_MAP_FORCE_DEVICEPTR:
> + assert (k->host_end - k->host_start == sizeof (void *));
> +
> + devicep->host2dev_func (devicep->target_id,
> + (void *) (tgt->tgt_start
> + + k->tgt_offset),
> + (void *) k->host_start,
> + sizeof (void *));
> + break;
> + default:
> + gomp_fatal ("%s: unhandled kind 0x%.2x", __FUNCTION__,
> + kind);
> }
> array++;
> }
> @@ -543,7 +540,7 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
> #undef GFC_DTYPE_TYPE_MASK
> #undef GFC_DTYPE_TYPE_SHIFT
> #undef GFC_DTYPE_SIZE_SHIFT
> -
> +
> if (is_target)
> {
> for (i = 0; i < mapnum; i++)
> @@ -1232,7 +1229,7 @@ gomp_target_init (void)
> 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 & TARGET_CAP_OPENACC_200)
> - ACC_register (&devices[i]);
> + goacc_register (&devices[i]);
> }
>
> free (offload_images);
> diff --git libgomp/task.c libgomp/task.c
> index 7d3233c..e9fa6ed 100644
> --- libgomp/task.c
> +++ libgomp/task.c
> @@ -1,7 +1,8 @@
> /* Copyright (C) 2007-2014 Free Software Foundation, Inc.
> Contributed by Richard Henderson <rth@redhat.com>.
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> diff --git libgomp/team.c libgomp/team.c
> index 594127c..24ddd07 100644
> --- libgomp/team.c
> +++ libgomp/team.c
> @@ -1,7 +1,8 @@
> /* Copyright (C) 2005-2014 Free Software Foundation, Inc.
> Contributed by Richard Henderson <rth@redhat.com>.
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> diff --git libgomp/testsuite/Makefile.in libgomp/testsuite/Makefile.in
> index c176f3a..78b6351 100644
> --- libgomp/testsuite/Makefile.in
> +++ libgomp/testsuite/Makefile.in
> [...]
> diff --git libgomp/testsuite/lib/libgomp.exp libgomp/testsuite/lib/libgomp.exp
> index 5f0755e..9ee71c4 100644
> --- libgomp/testsuite/lib/libgomp.exp
> +++ libgomp/testsuite/lib/libgomp.exp
> @@ -33,7 +33,7 @@ load_gcc_lib torture-options.exp
> load_gcc_lib fortran-modules.exp
>
> # Try to load a test support file, built during libgomp configuration.
> -load_file ../plugin/libgomp-test-support.exp
> +load_file libgomp-test-support.exp
>
> set dg-do-what-default run
>
> @@ -109,7 +109,7 @@ proc libgomp_init { args } {
> }
>
> # Compute what needs to be put into LD_LIBRARY_PATH
> - set always_ld_library_path ".:${blddir}/.libs:${blddir}/plugin/.libs"
> + set always_ld_library_path ".:${blddir}/.libs"
>
> # Get offload-related variables from environment (exported by Makefile)
> set offload_targets [getenv OFFLOAD_TARGETS]
> diff --git libgomp/plugin/libgomp-test-support.exp.in libgomp/testsuite/libgomp-test-support.exp.in
> similarity index 100%
> rename from libgomp/plugin/libgomp-test-support.exp.in
> rename to libgomp/testsuite/libgomp-test-support.exp.in
> diff --git libgomp/testsuite/libgomp.oacc-c++/c++.exp libgomp/testsuite/libgomp.oacc-c++/c++.exp
> index 9d5bf0b..c1f1f83 100644
> --- libgomp/testsuite/libgomp.oacc-c++/c++.exp
> +++ libgomp/testsuite/libgomp.oacc-c++/c++.exp
> @@ -24,7 +24,7 @@ dg-init
> # XXX (TEMPORARY): Remove the -flto once that's properly integrated.
> lappend ALWAYS_CFLAGS "additional_flags=-fopenacc -flto"
>
> -# TODO. Switch into C++ mode. Otherwise, the libgomp.oacc-c-c++-common/*.c
> +# Switch into C++ mode. Otherwise, the libgomp.oacc-c-c++-common/*.c
> # files would be compiled as C files.
> set SAVE_GCC_UNDER_TEST "$GCC_UNDER_TEST"
> set GCC_UNDER_TEST "$GCC_UNDER_TEST -x c++"
> @@ -113,7 +113,7 @@ if { $lang_test_file_found } {
> }
> }
>
> -# TODO. See above.
> +# See above.
> set GCC_UNDER_TEST "$SAVE_GCC_UNDER_TEST"
>
> # All done.
> diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/if-1.c libgomp/testsuite/libgomp.oacc-c-c++-common/if-1.c
> index e289f40..2a77936 100644
> --- libgomp/testsuite/libgomp.oacc-c-c++-common/if-1.c
> +++ libgomp/testsuite/libgomp.oacc-c-c++-common/if-1.c
> @@ -540,8 +540,5 @@ main(int argc, char **argv)
> abort ();
> }
>
> -#ifdef XXX_TODO_ENTER_END_DATA
> -#endif
> -
> return 0;
> }
> diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/lib-11.c libgomp/testsuite/libgomp.oacc-c-c++-common/lib-11.c
> index b4583ae..eccdb8c 100644
> --- libgomp/testsuite/libgomp.oacc-c-c++-common/lib-11.c
> +++ libgomp/testsuite/libgomp.oacc-c-c++-common/lib-11.c
> @@ -1,4 +1,5 @@
> -/* { dg-do run } */
> +/* Only nvptx plugin does the required error checking.
> + { dg-do run { target openacc_nvidia_accel_selected } } */
>
> #include <stdlib.h>
> #include <openacc.h>
> diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/lib-38.c libgomp/testsuite/libgomp.oacc-c-c++-common/lib-38.c
> index 1e16a1d..05d8498 100644
> --- libgomp/testsuite/libgomp.oacc-c-c++-common/lib-38.c
> +++ libgomp/testsuite/libgomp.oacc-c-c++-common/lib-38.c
> @@ -50,9 +50,6 @@ main (int argc, char **argv)
> if (!d2)
> abort ();
>
> - if (d1 != d2)
> - abort ();
> -
> acc_copyout (h, N);
>
> for (i = 0; i < N; i++)
> diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/lib-9.c libgomp/testsuite/libgomp.oacc-c-c++-common/lib-9.c
> index a4cf7f2..84045db 100644
> --- libgomp/testsuite/libgomp.oacc-c-c++-common/lib-9.c
> +++ libgomp/testsuite/libgomp.oacc-c-c++-common/lib-9.c
> @@ -58,7 +58,7 @@ main (int argc, char **argv)
> acc_set_device_num (1, (acc_device_t) 0);
>
> devnum = acc_get_device_num (devtype);
> - if (devnum != 1)
> + if (devnum != 0)
> abort ();
> }
>
> diff --git libgomp/work.c libgomp/work.c
> index bb563b8..8195a60 100644
> --- libgomp/work.c
> +++ libgomp/work.c
> @@ -1,7 +1,8 @@
> /* Copyright (C) 2005-2014 Free Software Foundation, Inc.
> Contributed by Richard Henderson <rth@redhat.com>.
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> diff --git liboffloadmic/plugin/Makefile.am liboffloadmic/plugin/Makefile.am
> index 0baf70d..a814f0c 100644
> --- liboffloadmic/plugin/Makefile.am
> +++ liboffloadmic/plugin/Makefile.am
> @@ -5,7 +5,8 @@
> # Contributed by Ilya Verbin <ilya.verbin@intel.com> and
> # Andrey Turetskiy <andrey.turetskiy@intel.com>.
> #
> -# This file is part of the GNU OpenMP Library (libgomp).
> +# 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
> diff --git liboffloadmic/plugin/Makefile.in liboffloadmic/plugin/Makefile.in
> index 5ba750a..21ce060 100644
> --- liboffloadmic/plugin/Makefile.in
> +++ liboffloadmic/plugin/Makefile.in
> [...]
> diff --git liboffloadmic/plugin/configure.ac liboffloadmic/plugin/configure.ac
> index 283faad..a2dd02d 100644
> --- liboffloadmic/plugin/configure.ac
> +++ liboffloadmic/plugin/configure.ac
> @@ -4,7 +4,8 @@
> #
> # Contributed by Andrey Turetskiy <andrey.turetskiy@intel.com>.
> #
> -# This file is part of the GNU OpenMP Library (libgomp).
> +# 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
> diff --git liboffloadmic/plugin/libgomp-plugin-intelmic.cpp liboffloadmic/plugin/libgomp-plugin-intelmic.cpp
> index 28ddbc3..0428b79 100644
> --- liboffloadmic/plugin/libgomp-plugin-intelmic.cpp
> +++ liboffloadmic/plugin/libgomp-plugin-intelmic.cpp
> @@ -4,7 +4,8 @@
>
> Contributed by Ilya Verbin <ilya.verbin@intel.com>.
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> diff --git liboffloadmic/plugin/offload_target_main.cpp liboffloadmic/plugin/offload_target_main.cpp
> index 4a2778e..3fead01 100644
> --- liboffloadmic/plugin/offload_target_main.cpp
> +++ liboffloadmic/plugin/offload_target_main.cpp
> @@ -4,7 +4,8 @@
>
> Contributed by Ilya Verbin <ilya.verbin@intel.com>.
>
> - This file is part of the GNU OpenMP Library (libgomp).
> + 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
> diff --git libstdc++-v3/doc/xml/manual/parallel_mode.xml libstdc++-v3/doc/xml/manual/parallel_mode.xml
> index 8ddec65..abf63ca 100644
> --- libstdc++-v3/doc/xml/manual/parallel_mode.xml
> +++ libstdc++-v3/doc/xml/manual/parallel_mode.xml
> @@ -106,7 +106,9 @@ It might work with other compilers, though.</para>
> not difficult: just compile your application with the compiler
> flag <literal>-fopenmp</literal>. This will link
> in <code>libgomp</code>, the
> - OpenMP <link xmlns:xlink="http://www.w3.org/1999/xlink"; xlink:href="http://gcc.gnu.org/onlinedocs/libgomp/";>GNU implementation</link>,
> + <link xmlns:xlink="http://www.w3.org/1999/xlink";
> + xlink:href="http://gcc.gnu.org/onlinedocs/libgomp/";>GNU Offloading and
> + Multi Processing Runtime Library</link>,
> whose presence is mandatory.
> </para>
>
GrÃÃe,
Thomas
Attachment:
pgpTbIfTMM6O0.pgp
Description: PGP signature
| Index Nav: | [Date Index] [Subject Index] [Author Index] [Thread Index] | |
|---|---|---|
| Message Nav: | [Date Prev] [Date Next] | [Thread Prev] [Thread Next] |