Re: [PATCH 3/4] Add libgomp plugin for Intel MIC

Re: [PATCH 3/4] Add libgomp plugin for Intel MIC

[Aleksander Ivanyushenko]

On Wed, Nov 11, 2015 at 17:56:15 +0300, Aleksander Ivanyushenko wrote:
> On Mon, Aug 24, 2015 at 10:45:03 +0200, Jakub Jelinek wrote:
> > On Thu, Aug 06, 2015 at 05:34:56PM +0300, Maxim Blumental wrote:
> > >  Applied the idea with python script alternative. Review, please.
> > 
> > > 2015-07-28  Maxim Blumenthal  <maxim.blumenthal@intel.com>
> > > 
> > > 	* configure.ac: Add a check for xxd or python presence when the target
> > > 	is intelmic or intelmicemul.
> > > 	* configure: Regenerate.
> > > 	* liboffloadmic/plugin/Makefile.am: Add a condition into
> > > 	make_target_image.h generating code.  This condition performs an
> > > 	action with either xxd or a special python script during the
> > > 	generating.
> > > 	* liboffloadmic/plugin/xxd.py: New file.
> > > 	* liboffloadmic/plugin/Makefile.in: Regenerate.
> > 
> > I still don't like this, there should be no `which ...` uses in the
> > Makefile.
> > Instead, use AC_CHECK_PROG/AC_CHECK_PROGS in configure.ac, for python
> > perhaps search for python python2 python3 or what is common in the python
> > land.  And prepare the command line to use in the Makefile.am in configure
> > too, then AC_SUBST it and use the variable in there (and the variable will
> > use $@ etc.).
> Maxim has left Intel so I have fixed this issue. I tried to build with and
> without xxd, so everything works fine. ok for trunk?
> 
> 2015-11-10  Aleksander Ivanushenko  <aleksander.ivanushenko@intel.com>
> 	    Maxim Blumenthal  <maxim.blumenthal@intel.com>
> 
> 	* configure.ac: Add xxd and python check for intelmic and
> 	intelmicemul.
> 	* configure: Regenerate.
> 
> liboffloadmic/
> 2015-11-10  Aleksander Ivanushenko  <aleksander.ivanushenko@intel.com>
> 	    Maxim Blumenthal  <maxim.blumenthal@intel.com>
> 	    David Malcolm  <dmalcolm@redhat.com>
> 
> 	* plugin/xxd.py: New file.
> 	* plugin/configure.ac: Add searching for xxd and python pathes.
> 	* plugin/Makefile.am: Add python script usage in case when xxd is not
> 	available.
> 	* plugin/configure: Regenerate.
> 	* plugin/Makefile.in: Regenerate.
> 
>
Ping. 

Re: [PATCH 3/4] Add libgomp plugin for Intel MIC

[Jakub Jelinek]

On Wed, Nov 11, 2015 at 05:56:15PM +0300, Aleksander Ivanyushenko wrote:
> diff --git a/configure.ac b/configure.ac
> index 9241261..b997646 100644
> --- a/configure.ac
> +++ b/configure.ac
> @@ -494,6 +494,18 @@ else
>  fi])
>  AC_SUBST(extra_liboffloadmic_configure_flags)
>  
> +# Intelmic and intelmicemul require xxd or python.
> +case "${target}" in
> +  *-intelmic-* | *-intelmicemul-*)
> +    AC_CHECK_PROG(xxd_present, xxd, "yes", "no")
> +    AC_CHECK_PROG(python2_present, python2, "yes", "no")
> +    AC_CHECK_PROG(python3_present, python3, "yes", "no")
> +    if test "$xxd_present$python2_present$python3_present" = "nonono"; then
> +      AC_MSG_ERROR([cannot find neither xxd nor python])
> +    fi
> +    ;;
> +esac

Why here?  I'd do something like that only in
liboffloadmic/plugin/configure.ac.  Furthermore, it is inconsistent
with what you actually use in liboffloadmic/plugin (where you look only
for python and above you only look for python[23]).

> @@ -73,7 +75,7 @@ main_target_image.h: offload_target_main
>  	@echo "};" >> $@
>  	@echo "extern \"C\" const MainTargetImage main_target_image = {" >> $@
>  	@echo "  image_size, \"offload_target_main\"," >> $@
> -	@cat $< | xxd -include >> $@
> +	@if test "x$(xxd_path)" != "xno"; then cat $< | $(xxd_path) -include >> $@; else $(python_path) $(XXD_PY) $< >> $@; fi;
>  	@echo "};" >> $@

I'd prefer to use $(XXD) and $(PYTHON) instead of $(xxd_path) and $(python_path),
that is more consistent with dozens of other variables for other tools.

> --- a/liboffloadmic/plugin/configure.ac
> +++ b/liboffloadmic/plugin/configure.ac
> @@ -124,6 +124,10 @@ case ${enable_version_specific_runtime_libs} in
>      ;;
>  esac
>  
> +# Find path to xxd or python
> +AC_PATH_PROG(xxd_path, xxd, "no")
> +AC_PATH_PROG(python_path, python, "no")

I'd use
+AC_PATH_PROG(XXD, xxd, no)
+AC_PATH_PROGS(PYTHON, python python2 python3, no)
and then add the conditional AC_MSG_ERROR if
x$XXD = xno && x$PYTHON = xno

	Jakub

[PATCH 0/4] OpenMP 4.0 offloading to Intel MIC

[Ilya Verbin]

Hello,

This patchset would contain target-specific things to support offloading to the
devices with Intel MIC architecture.
Particularly: mkoffload tool, liboffloadmic library and a plugin for libgomp.

  -- Ilya

[PATCH 3/4] Add libgomp plugin for Intel MIC

[Ilya Verbin]

Hello,

This patch contains a plugin for libgomp and appropriate changes for makefiles.

The plugin uses liboffloadmic_host.so to interact with the device (or with an
emulator).  Also the patch contains offload_target_main executable, which is the
corresponding target side part of a libgomp plugin, and it uses
liboffloadmic_target.so.

The plugin builds automatically with liboffloadmic.

Autogenerated files are skipped.  Is it ok for trunk?

Thanks,
  -- Ilya


2014-10-21  Ilya Verbin  <ilya.verbin@intel.com>
	    Andrey Turetskiy  <andrey.turetskiy@intel.com>

liboffloadmic/
	* Makefile.in: Regenerate.
	* configure: Regenerate.
	* configure.ac: Add subdirectory 'plugin'.
	* plugin/Makefile.am: New file.
	* plugin/Makefile.in: New file, generated by automake.
	* plugin/aclocal.m4: New file, generated by aclocal.
	* plugin/configure: New file, generated by autoconf.
	* plugin/configure.ac: New file.
	* plugin/libgomp-plugin-intelmic.cpp: New file.
	* plugin/offload_target_main.cpp: New file.

---

diff --git a/liboffloadmic/configure.ac b/liboffloadmic/configure.ac
index fb575b3..81fae8f 100644
--- a/liboffloadmic/configure.ac
+++ b/liboffloadmic/configure.ac
@@ -42,6 +42,7 @@ AC_PROG_CC
 AC_PROG_CXX
 AC_CONFIG_FILES([Makefile liboffloadmic_host.spec liboffloadmic_target.spec])
 AM_ENABLE_MULTILIB(, ..)
+AC_CONFIG_SUBDIRS(plugin)
 AC_FUNC_ALLOCA
 AC_CHECK_HEADERS([mm_malloc.h], [], [AC_MSG_ERROR(["Couldn't find mm_malloc.h"])])
 AC_CHECK_FUNCS([__secure_getenv secure_getenv])
diff --git a/liboffloadmic/plugin/Makefile.am b/liboffloadmic/plugin/Makefile.am
new file mode 100644
index 0000000..0baf70d
--- /dev/null
+++ b/liboffloadmic/plugin/Makefile.am
@@ -0,0 +1,123 @@
+# Plugin for offload execution on Intel MIC devices.
+#
+# Copyright (C) 2014 Free Software Foundation, Inc.
+#
+# 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).
+#
+# 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/>.
+
+
+AUTOMAKE_OPTIONS = foreign
+ACLOCAL_AMFLAGS = -I ../.. -I ../../config
+
+# Directories
+build_dir = $(top_builddir)
+source_dir = $(top_srcdir)
+coi_inc_dir = $(top_srcdir)/../include/coi
+myo_inc_dir = $(top_srcdir)/../include/myo
+libgomp_src_dir = $(top_srcdir)/../../libgomp
+libgomp_dir = $(build_dir)/../../libgomp
+liboffload_src_dir = $(top_srcdir)/../runtime
+liboffload_dir = $(top_builddir)/..
+
+# May be used by toolexeclibdir.
+gcc_version := $(shell cat $(top_srcdir)/../../gcc/BASE-VER)
+libsubincludedir = $(libdir)/gcc/$(target_alias)/$(gcc_version)$(MULTISUBDIR)/include
+# Search for main_target_image.h in these directories
+target_prefix_dir = $(libdir)/gcc/$(accel_target)/$(gcc_version)$(MULTISUBDIR)
+target_build_dir = $(accel_search_dir)/$(accel_target)$(MULTISUBDIR)/liboffloadmic/plugin
+target_install_dir = $(accel_search_dir)/lib/gcc/$(accel_target)/$(gcc_version)$(MULTISUBDIR)
+
+if PLUGIN_HOST
+  toolexeclib_LTLIBRARIES = libgomp-plugin-intelmic.la
+  libgomp_plugin_intelmic_la_SOURCES = libgomp-plugin-intelmic.cpp
+  libgomp_plugin_intelmic_la_CPPFLAGS = $(CPPFLAGS) -DLINUX -DCOI_LIBRARY_VERSION=2 -DMYO_SUPPORT -DOFFLOAD_DEBUG=1 -DSEP_SUPPORT -DTIMING_SUPPORT -DHOST_LIBRARY=1 -I$(coi_inc_dir) -I$(myo_inc_dir) -I$(liboffload_src_dir) -I$(libgomp_src_dir) -I$(libgomp_dir) -I$(target_prefix_dir)/include -I$(target_build_dir) -I$(target_install_dir)/include
+  libgomp_plugin_intelmic_la_LDFLAGS = -L$(liboffload_dir)/.libs -loffloadmic_host -version-info 1:0:0
+else # PLUGIN_TARGET
+  plugin_includedir = $(libsubincludedir)
+  plugin_include_HEADERS = main_target_image.h
+  AM_CPPFLAGS = $(CPPFLAGS) -DLINUX -DCOI_LIBRARY_VERSION=2 -DMYO_SUPPORT -DOFFLOAD_DEBUG=1 -DSEP_SUPPORT -DTIMING_SUPPORT -DHOST_LIBRARY=0 -I$(coi_inc_dir) -I$(myo_inc_dir) -I$(liboffload_src_dir) -I$(libgomp_dir)
+  AM_CXXFLAGS = $(CXXFLAGS)
+  AM_LDFLAGS = -L$(liboffload_dir)/.libs -L$(libgomp_dir)/.libs -loffloadmic_target -lcoi_device -lmyo-service -lgomp -rdynamic
+endif
+
+main_target_image.h: offload_target_main
+	@echo -n "const int image_size = " > $@
+	@stat -c '%s' $< >> $@
+	@echo ";" >> $@
+	@echo "struct MainTargetImage {" >> $@
+	@echo "  int64_t size;" >> $@
+	@echo "  char name[sizeof \"offload_target_main\"];" >> $@
+	@echo "  char data[image_size];" >> $@
+	@echo "};" >> $@
+	@echo "extern \"C\" const MainTargetImage main_target_image = {" >> $@
+	@echo "  image_size, \"offload_target_main\"," >> $@
+	@cat $< | xxd -include >> $@
+	@echo "};" >> $@
+
+offload_target_main: $(liboffload_dir)/ofldbegin.o offload_target_main.o $(liboffload_dir)/ofldend.o
+	$(CXX) $(AM_LDFLAGS) $^ -o $@
+
+offload_target_main.o: offload_target_main.cpp
+	$(CXX) $(AM_CXXFLAGS) $(AM_CPPFLAGS) -c $< -o $@
+
+# Work around what appears to be a GNU make bug handling MAKEFLAGS
+# values defined in terms of make variables, as is the case for CC and
+# friends when we are called from the top level Makefile.
+AM_MAKEFLAGS = \
+       "AR_FLAGS=$(AR_FLAGS)" \
+       "CC_FOR_BUILD=$(CC_FOR_BUILD)" \
+       "CFLAGS=$(CFLAGS)" \
+       "CXXFLAGS=$(CXXFLAGS)" \
+       "CFLAGS_FOR_BUILD=$(CFLAGS_FOR_BUILD)" \
+       "CFLAGS_FOR_TARGET=$(CFLAGS_FOR_TARGET)" \
+       "INSTALL=$(INSTALL)" \
+       "INSTALL_DATA=$(INSTALL_DATA)" \
+       "INSTALL_PROGRAM=$(INSTALL_PROGRAM)" \
+       "INSTALL_SCRIPT=$(INSTALL_SCRIPT)" \
+       "JC1FLAGS=$(JC1FLAGS)" \
+       "LDFLAGS=$(LDFLAGS)" \
+       "LIBCFLAGS=$(LIBCFLAGS)" \
+       "LIBCFLAGS_FOR_TARGET=$(LIBCFLAGS_FOR_TARGET)" \
+       "MAKE=$(MAKE)" \
+       "MAKEINFO=$(MAKEINFO) $(MAKEINFOFLAGS)" \
+       "PICFLAG=$(PICFLAG)" \
+       "PICFLAG_FOR_TARGET=$(PICFLAG_FOR_TARGET)" \
+       "SHELL=$(SHELL)" \
+       "RUNTESTFLAGS=$(RUNTESTFLAGS)" \
+       "exec_prefix=$(exec_prefix)" \
+       "infodir=$(infodir)" \
+       "libdir=$(libdir)" \
+       "prefix=$(prefix)" \
+       "includedir=$(includedir)" \
+       "AR=$(AR)" \
+       "AS=$(AS)" \
+       "LD=$(LD)" \
+       "LIBCFLAGS=$(LIBCFLAGS)" \
+       "NM=$(NM)" \
+       "PICFLAG=$(PICFLAG)" \
+       "RANLIB=$(RANLIB)" \
+       "DESTDIR=$(DESTDIR)"
+
+MAKEOVERRIDES =
+
diff --git a/liboffloadmic/plugin/configure.ac b/liboffloadmic/plugin/configure.ac
new file mode 100644
index 0000000..283faad
--- /dev/null
+++ b/liboffloadmic/plugin/configure.ac
@@ -0,0 +1,135 @@
+# Plugin for offload execution on Intel MIC devices.
+#
+# Copyright (C) 2014 Free Software Foundation, Inc.
+#
+# Contributed by Andrey Turetskiy <andrey.turetskiy@intel.com>.
+#
+# This file is part of the GNU OpenMP 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/>.
+
+# Process this file with autoconf to produce a configure script, like so:
+# aclocal -I ../.. -I ../../config && autoconf && automake
+
+AC_PREREQ([2.64])
+AC_INIT([Intel MIC Offload Plugin], [1.0], ,[libgomp-plugin-intelmic])
+
+AC_CONFIG_AUX_DIR(../..)
+
+AC_CANONICAL_SYSTEM
+target_alias=${target_alias-$host_alias}
+AC_SUBST(target_alias)
+
+AM_INIT_AUTOMAKE([1.9.0 foreign no-dist])
+
+AM_MAINTAINER_MODE
+
+AC_PROG_CC
+AC_PROG_CXX
+AC_CONFIG_FILES([Makefile])
+AM_ENABLE_MULTILIB(, ../..)
+
+if test "${multilib}" = "yes"; then
+  multilib_arg="--enable-multilib"
+else
+  multilib_arg=
+fi
+
+# Make sure liboffloadmic is enabled
+case "$enable_liboffloadmic" in
+  host | target)
+    ;;
+  *)
+    AC_MSG_ERROR([Liboffloadmic is disabled]) ;;
+esac
+AM_CONDITIONAL(PLUGIN_HOST, [test x"$enable_liboffloadmic" = xhost])
+
+# Get accel target and path to build or install tree of accel compiler
+accel_search_dir=
+accel_target=
+if test x"$enable_liboffloadmic" = xhost; then
+  for accel in `echo $enable_offload_targets | sed -e 's#,# #g'`; do
+    accel_name=`echo $accel | sed 's/=.*//'`
+    accel_dir=`echo $accel | grep '=' | sed 's/.*=//'`
+    case "$accel_name" in
+      *-intelmic-* | *-intelmicemul-*)
+	accel_target=$accel_name
+	accel_search_dir=$accel_dir
+	;;
+    esac
+  done
+  if test x"$accel_target" = x; then
+    AC_MSG_ERROR([--enable-offload-targets does not contain intelmic target])
+  fi
+fi
+AC_SUBST(accel_search_dir)
+AC_SUBST(accel_target)
+
+AC_MSG_CHECKING([for --enable-version-specific-runtime-libs])
+AC_ARG_ENABLE([version-specific-runtime-libs],
+  AC_HELP_STRING([--enable-version-specific-runtime-libs],
+		 [Specify that runtime libraries should be installed in a compiler-specific directory]),
+  [case "$enableval" in
+    yes) enable_version_specific_runtime_libs=yes ;;
+    no)  enable_version_specific_runtime_libs=no ;;
+    *)   AC_MSG_ERROR([Unknown argument to enable/disable version-specific libs]);;
+   esac],
+  [enable_version_specific_runtime_libs=no])
+AC_MSG_RESULT($enable_version_specific_runtime_libs)
+
+
+# Calculate toolexeclibdir.
+# Also toolexecdir, though it's only used in toolexeclibdir.
+case ${enable_version_specific_runtime_libs} in
+  yes)
+    # Need the gcc compiler version to know where to install libraries
+    # and header files if --enable-version-specific-runtime-libs option
+    # is selected.
+    toolexecdir='$(libdir)/gcc/$(target_alias)'
+    toolexeclibdir='$(toolexecdir)/$(gcc_version)$(MULTISUBDIR)'
+    ;;
+  no)
+    if test -n "$with_cross_host" &&
+       test x"$with_cross_host" != x"no"; then
+      # Install a library built with a cross compiler in tooldir, not libdir.
+      toolexecdir='$(exec_prefix)/$(target_alias)'
+      toolexeclibdir='$(toolexecdir)/lib'
+    else
+      toolexecdir='$(libdir)/gcc-lib/$(target_alias)'
+      toolexeclibdir='$(libdir)'
+    fi
+    multi_os_directory=`$CC -print-multi-os-directory`
+    case $multi_os_directory in
+      .) ;; # Avoid trailing /.
+      *) toolexeclibdir=$toolexeclibdir/$multi_os_directory ;;
+    esac
+    ;;
+esac
+
+AC_LIBTOOL_DLOPEN
+AM_PROG_LIBTOOL
+# Forbid libtool to hardcode RPATH, because we want to be able to specify
+# library search directory using LD_LIBRARY_PATH
+hardcode_into_libs=no
+AC_SUBST(toolexecdir)
+AC_SUBST(toolexeclibdir)
+
+# Must be last
+AC_OUTPUT
diff --git a/liboffloadmic/plugin/libgomp-plugin-intelmic.cpp b/liboffloadmic/plugin/libgomp-plugin-intelmic.cpp
new file mode 100644
index 0000000..dbbeeaf
--- /dev/null
+++ b/liboffloadmic/plugin/libgomp-plugin-intelmic.cpp
@@ -0,0 +1,442 @@
+/* Plugin for offload execution on Intel MIC devices.
+
+   Copyright (C) 2014 Free Software Foundation, Inc.
+
+   Contributed by Ilya Verbin <ilya.verbin@intel.com>.
+
+   This file is part of the GNU OpenMP 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/>.  */
+
+/* Host side part of a libgomp plugin.  */
+
+#include <stdint.h>
+#include <stdio.h>
+#include <stdlib.h>
+#include <string.h>
+#include <utility>
+#include <vector>
+#include <libgomp_target.h>
+#include "compiler_if_host.h"
+#include "main_target_image.h"
+
+#define LD_LIBRARY_PATH_ENV	"LD_LIBRARY_PATH"
+#define MIC_LD_LIBRARY_PATH_ENV	"MIC_LD_LIBRARY_PATH"
+
+#ifdef DEBUG
+#define TRACE(...)					    \
+{							    \
+fprintf (stderr, "HOST:\t%s:%s ", __FILE__, __FUNCTION__);  \
+fprintf (stderr, __VA_ARGS__);				    \
+fprintf (stderr, "\n");					    \
+}
+#else
+#define TRACE { }
+#endif
+
+
+static VarDesc vd_host2tgt = {
+  { 1, 1 },		      /* dst, src			      */
+  { 1, 0 },		      /* in, out			      */
+  1,			      /* alloc_if			      */
+  1,			      /* free_if			      */
+  4,			      /* align				      */
+  0,			      /* mic_offset			      */
+  { 0, 0, 0, 0, 0, 0, 0, 0 }, /* is_static, is_static_dstn, has_length,
+				 is_stack_buf, sink_addr, alloc_disp,
+				 is_noncont_src, is_noncont_dst	      */
+  0,			      /* offset				      */
+  0,			      /* size				      */
+  1,			      /* count				      */
+  0,			      /* alloc				      */
+  0,			      /* into				      */
+  0			      /* ptr				      */
+};
+
+static VarDesc vd_tgt2host = {
+  { 1, 1 },		      /* dst, src			      */
+  { 0, 1 },		      /* in, out			      */
+  1,			      /* alloc_if			      */
+  1,			      /* free_if			      */
+  4,			      /* align				      */
+  0,			      /* mic_offset			      */
+  { 0, 0, 0, 0, 0, 0, 0, 0 }, /* is_static, is_static_dstn, has_length,
+				 is_stack_buf, sink_addr, alloc_disp,
+				 is_noncont_src, is_noncont_dst	      */
+  0,			      /* offset				      */
+  0,			      /* size				      */
+  1,			      /* count				      */
+  0,			      /* alloc				      */
+  0,			      /* into				      */
+  0			      /* ptr				      */
+};
+
+
+/* Total number of shared libraries with offloading to Intel MIC.  */
+static int num_libraries;
+
+/* Pointers to the descriptors, containing pointers to host-side tables and to
+   target images.  */
+static std::vector< std::pair<void *, void *> > lib_descrs;
+
+/* Mutex to control parallel plugin calls.  */
+static pthread_mutex_t mutex = PTHREAD_MUTEX_INITIALIZER;
+
+
+/* Add path specified in LD_LIBRARY_PATH to MIC_LD_LIBRARY_PATH, which is
+   required by liboffloadmic.  */
+__attribute__((constructor))
+static void
+set_mic_lib_path (void)
+{
+  const char *ld_lib_path = getenv (LD_LIBRARY_PATH_ENV);
+  const char *mic_lib_path = getenv (MIC_LD_LIBRARY_PATH_ENV);
+  char *mic_lib_path_new;
+
+  if (!ld_lib_path)
+    return;
+
+  mic_lib_path_new = (char *) malloc ((mic_lib_path ? strlen (mic_lib_path) : 0)
+				      + strlen (ld_lib_path) + 2);
+
+  if (!mic_lib_path)
+    strcpy (mic_lib_path_new, ld_lib_path);
+  else
+    sprintf (mic_lib_path_new, "%s:%s", mic_lib_path, ld_lib_path);
+  setenv (MIC_LD_LIBRARY_PATH_ENV, mic_lib_path_new, 1);
+  free (mic_lib_path_new);
+}
+
+extern "C" enum offload_target_type
+GOMP_OFFLOAD_get_type (void)
+{
+  enum offload_target_type res = OFFLOAD_TARGET_TYPE_INTEL_MIC;
+  TRACE ("(): return %d", res);
+  return res;
+}
+
+extern "C" int
+GOMP_OFFLOAD_get_num_devices (void)
+{
+  int res = _Offload_number_of_devices ();
+  TRACE ("(): return %d", res);
+  return res;
+}
+
+/* This should be called from every shared library with offloading.  */
+extern "C" void
+GOMP_OFFLOAD_register_image (void *host_table, void *target_image)
+{
+  TRACE ("(host_table = %p, target_image = %p)", host_table, target_image);
+
+  if (num_libraries >= 1000)
+    {
+      fprintf (stderr,
+	       "%s: The number of loaded shared libraries is over 1000!\n",
+	       __FILE__);
+      exit (1);
+    }
+
+  lib_descrs.push_back (std::make_pair (host_table, target_image));
+  num_libraries++;
+}
+
+static void
+offload (const char *file, uint64_t line, int device, const char *name,
+	 int num_vars, VarDesc *vars, VarDesc2 *vars2)
+{
+  OFFLOAD ofld = __offload_target_acquire1 (&device, file, line);
+  if (ofld)
+    __offload_offload1 (ofld, name, 0, num_vars, vars, vars2, 0, NULL, NULL);
+  else
+    {
+      fprintf (stderr, "%s:%d: Offload target acquire failed\n", file, line);
+      exit (1);
+    }
+}
+
+static int first_init = 1;
+
+/* Load offload_target_main on target.  */
+extern "C" void
+GOMP_OFFLOAD_init_device (int device)
+{
+  TRACE ("");
+  pthread_mutex_lock (&mutex);
+  if (first_init)
+    {
+      __offload_register_image (&main_target_image);
+      first_init = 0;
+    }
+  pthread_mutex_unlock (&mutex);
+  offload (__FILE__, __LINE__, device, "__offload_target_init_proc", 0,
+	   NULL, NULL);
+}
+
+static void
+get_target_table (int device, int &num_funcs, int &num_vars, void **&table)
+{
+  VarDesc vd1[2] = { vd_tgt2host, vd_tgt2host };
+  vd1[0].ptr = &num_funcs;
+  vd1[0].size = sizeof (num_funcs);
+  vd1[1].ptr = &num_vars;
+  vd1[1].size = sizeof (num_vars);
+  VarDesc2 vd1g[2] = { { "num_funcs", 0 }, { "num_vars", 0 } };
+
+  offload (__FILE__, __LINE__, device, "__offload_target_table_p1", 2,
+	   vd1, vd1g);
+
+  int table_size = num_funcs + 2 * num_vars;
+  if (table_size > 0)
+    {
+      table = new void * [table_size];
+
+      VarDesc vd2;
+      vd2 = vd_tgt2host;
+      vd2.ptr = table;
+      vd2.size = table_size * sizeof (void *);
+      VarDesc2 vd2g = { "table", 0 };
+
+      offload (__FILE__, __LINE__, device, "__offload_target_table_p2", 1,
+	       &vd2, &vd2g);
+    }
+}
+
+static void
+load_lib_and_get_table (int device, int lib_num, mapping_table *&table,
+			int &table_size)
+{
+  /* 1000 shared libraries with offloading ought to be enough for anybody.  */
+  struct TargetImage {
+    int64_t size;
+    char name[sizeof "lib000.so"];
+    char data[];
+  } __attribute__ ((packed));
+
+  void ***host_table_descr = (void ***) lib_descrs[lib_num].first;
+  void **host_func_start = host_table_descr[0];
+  void **host_func_end   = host_table_descr[1];
+  void **host_var_start  = host_table_descr[2];
+  void **host_var_end    = host_table_descr[3];
+
+  void **target_image_descr = (void **) lib_descrs[lib_num].second;
+  void *image_start = target_image_descr[0];
+  void *image_end   = target_image_descr[1];
+
+  TRACE ("() host_table_descr { %p, %p, %p, %p }", host_func_start,
+	 host_func_end, host_var_start, host_var_end);
+  TRACE ("() target_image_descr { %p, %p }", image_start, image_end);
+
+  int64_t image_size = (uintptr_t) image_end - (uintptr_t) image_start;
+  TargetImage *image
+    = (TargetImage *) malloc (sizeof (int64_t) + sizeof ("lib000.so")
+			      + image_size);
+  image->size = image_size;
+  sprintf (image->name, "lib%03d.so", lib_num);
+  memcpy (image->data, image_start, image->size);
+
+  TRACE ("() __offload_register_image %s { %p, %d }",
+	 image->name, image_start, image->size);
+  __offload_register_image (image);
+
+  int tgt_num_funcs = 0;
+  int tgt_num_vars = 0;
+  void **tgt_table = NULL;
+  get_target_table (device, tgt_num_funcs, tgt_num_vars, tgt_table);
+  free (image);
+
+  /* The func table contains only addresses, the var table contains addresses
+     and corresponding sizes.  */
+  int host_num_funcs = host_func_end - host_func_start;
+  int host_num_vars  = (host_var_end - host_var_start) / 2;
+  TRACE ("() host_num_funcs = %d, tgt_num_funcs = %d",
+	 host_num_funcs, tgt_num_funcs);
+  TRACE ("() host_num_vars = %d, tgt_num_vars = %d",
+	 host_num_vars, tgt_num_vars);
+  if (host_num_funcs != tgt_num_funcs)
+    {
+      fprintf (stderr, "%s: Can't map target functions\n", __FILE__);
+      exit (1);
+    }
+  if (host_num_vars != tgt_num_vars)
+    {
+      fprintf (stderr, "%s: Can't map target variables\n", __FILE__);
+      exit (1);
+    }
+
+  table = (mapping_table *) realloc (table, (table_size + host_num_funcs
+					     + host_num_vars)
+					    * sizeof (mapping_table));
+  if (table == NULL)
+    {
+      fprintf (stderr, "%s: Can't allocate memory\n", __FILE__);
+      exit (1);
+    }
+
+  for (int i = 0; i < host_num_funcs; i++)
+    {
+      mapping_table t;
+      t.host_start = (uintptr_t) host_func_start[i];
+      t.host_end = t.host_start + 1;
+      t.tgt_start = (uintptr_t) tgt_table[i];
+      t.tgt_end = t.tgt_start + 1;
+
+      TRACE ("() lib %d, func %d:\t0x%llx -- 0x%llx",
+	     lib_num, i, t.host_start, t.tgt_start);
+
+      table[table_size++] = t;
+    }
+
+  for (int i = 0; i < host_num_vars * 2; i += 2)
+    {
+      mapping_table t;
+      t.host_start = (uintptr_t) host_var_start[i];
+      t.host_end = t.host_start + (uintptr_t) host_var_start[i+1];
+      t.tgt_start = (uintptr_t) tgt_table[tgt_num_funcs+i];
+      t.tgt_end = t.tgt_start + (uintptr_t) tgt_table[tgt_num_funcs+i+1];
+
+      TRACE ("() lib %d, var %d:\t0x%llx (%d) -- 0x%llx (%d)", lib_num, i/2,
+	     t.host_start, t.host_end - t.host_start,
+	     t.tgt_start, t.tgt_end - t.tgt_start);
+
+      table[table_size++] = t;
+    }
+
+  delete [] tgt_table;
+}
+
+extern "C" int
+GOMP_OFFLOAD_get_table (int device, void *result)
+{
+  TRACE ("(num_libraries = %d)", num_libraries);
+
+  mapping_table *table = NULL;
+  int table_size = 0;
+
+  for (int i = 0; i < num_libraries; i++)
+    load_lib_and_get_table (device, i, table, table_size);
+
+  *(void **) result = table;
+  return table_size;
+}
+
+extern "C" void *
+GOMP_OFFLOAD_alloc (int device, size_t size)
+{
+  TRACE ("(size = %d)", size);
+
+  void *tgt_ptr;
+  VarDesc vd1[2] = { vd_host2tgt, vd_tgt2host };
+  vd1[0].ptr = &size;
+  vd1[0].size = sizeof (size);
+  vd1[1].ptr = &tgt_ptr;
+  vd1[1].size = sizeof (void *);
+  VarDesc2 vd1g[2] = { { "size", 0 }, { "tgt_ptr", 0 } };
+
+  offload (__FILE__, __LINE__, device, "__offload_target_alloc", 2, vd1, vd1g);
+
+  return tgt_ptr;
+}
+
+extern "C" void
+GOMP_OFFLOAD_free (int device, void *tgt_ptr)
+{
+  TRACE ("(tgt_ptr = %p)", tgt_ptr);
+
+  VarDesc vd1 = vd_host2tgt;
+  vd1.ptr = &tgt_ptr;
+  vd1.size = sizeof (void *);
+  VarDesc2 vd1g = { "tgt_ptr", 0 };
+
+  offload (__FILE__, __LINE__, device, "__offload_target_free", 1, &vd1, &vd1g);
+}
+
+extern "C" void *
+GOMP_OFFLOAD_host2dev (int device, void *tgt_ptr, const void *host_ptr,
+		       size_t size)
+{
+  TRACE ("(tgt_ptr = %p, host_ptr = %p, size = %d)", tgt_ptr, host_ptr, size);
+  if (!size)
+    return tgt_ptr;
+
+  VarDesc vd1[2] = { vd_host2tgt, vd_host2tgt };
+  vd1[0].ptr = &tgt_ptr;
+  vd1[0].size = sizeof (void *);
+  vd1[1].ptr = &size;
+  vd1[1].size = sizeof (size);
+  VarDesc2 vd1g[2] = { { "tgt_ptr", 0 }, { "size", 0 } };
+
+  offload (__FILE__, __LINE__, device, "__offload_target_host2tgt_p1", 2,
+	   vd1, vd1g);
+
+  VarDesc vd2 = vd_host2tgt;
+  vd2.ptr = (void *) host_ptr;
+  vd2.size = size;
+  VarDesc2 vd2g = { "var", 0 };
+
+  offload (__FILE__, __LINE__, device, "__offload_target_host2tgt_p2", 1,
+	   &vd2, &vd2g);
+
+  return tgt_ptr;
+}
+
+extern "C" void *
+GOMP_OFFLOAD_dev2host (int device, void *host_ptr, const void *tgt_ptr,
+		       size_t size)
+{
+  TRACE ("(host_ptr = %p, tgt_ptr = %p, size = %d)", host_ptr, tgt_ptr, size);
+  if (!size)
+    return host_ptr;
+
+  VarDesc vd1[2] = { vd_host2tgt, vd_host2tgt };
+  vd1[0].ptr = &tgt_ptr;
+  vd1[0].size = sizeof (void *);
+  vd1[1].ptr = &size;
+  vd1[1].size = sizeof (size);
+  VarDesc2 vd1g[2] = { { "tgt_ptr", 0 }, { "size", 0 } };
+
+  offload (__FILE__, __LINE__, device, "__offload_target_tgt2host_p1", 2,
+	   vd1, vd1g);
+
+  VarDesc vd2 = vd_tgt2host;
+  vd2.ptr = (void *) host_ptr;
+  vd2.size = size;
+  VarDesc2 vd2g = { "var", 0 };
+
+  offload (__FILE__, __LINE__, device, "__offload_target_tgt2host_p2", 1,
+	   &vd2, &vd2g);
+
+  return host_ptr;
+}
+
+extern "C" void
+GOMP_OFFLOAD_run (int device, void *tgt_fn, void *tgt_vars)
+{
+  TRACE ("(tgt_fn = %p, tgt_vars = %p)", tgt_fn, tgt_vars);
+
+  VarDesc vd1[2] = { vd_host2tgt, vd_host2tgt };
+  vd1[0].ptr = &tgt_fn;
+  vd1[0].size = sizeof (void *);
+  vd1[1].ptr = &tgt_vars;
+  vd1[1].size = sizeof (void *);
+  VarDesc2 vd1g[2] = { { "tgt_fn", 0 }, { "tgt_vars", 0 } };
+
+  offload (__FILE__, __LINE__, device, "__offload_target_run", 2, vd1, vd1g);
+}
diff --git a/liboffloadmic/plugin/offload_target_main.cpp b/liboffloadmic/plugin/offload_target_main.cpp
new file mode 100644
index 0000000..4a2778e
--- /dev/null
+++ b/liboffloadmic/plugin/offload_target_main.cpp
@@ -0,0 +1,366 @@
+/* Plugin for offload execution on Intel MIC devices.
+
+   Copyright (C) 2014 Free Software Foundation, Inc.
+
+   Contributed by Ilya Verbin <ilya.verbin@intel.com>.
+
+   This file is part of the GNU OpenMP 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/>.  */
+
+/* Target side part of a libgomp plugin.  */
+
+#include <stdint.h>
+#include <stdio.h>
+#include <stdlib.h>
+#include "compiler_if_target.h"
+
+
+#ifdef DEBUG
+#define TRACE(...)					      \
+{							      \
+fprintf (stderr, "TARGET:\t%s:%s ", __FILE__, __FUNCTION__);  \
+fprintf (stderr, __VA_ARGS__);				      \
+fprintf (stderr, "\n");					      \
+}
+#else
+#define TRACE { }
+#endif
+
+
+static VarDesc vd_host2tgt = {
+  { 1, 1 },		      /* dst, src			      */
+  { 1, 0 },		      /* in, out			      */
+  1,			      /* alloc_if			      */
+  1,			      /* free_if			      */
+  4,			      /* align				      */
+  0,			      /* mic_offset			      */
+  { 0, 0, 0, 0, 0, 0, 0, 0 }, /* is_static, is_static_dstn, has_length,
+				 is_stack_buf, sink_addr, alloc_disp,
+				 is_noncont_src, is_noncont_dst	      */
+  0,			      /* offset				      */
+  0,			      /* size				      */
+  1,			      /* count				      */
+  0,			      /* alloc				      */
+  0,			      /* into				      */
+  0			      /* ptr				      */
+};
+
+static VarDesc vd_tgt2host = {
+  { 1, 1 },		      /* dst, src			      */
+  { 0, 1 },		      /* in, out			      */
+  1,			      /* alloc_if			      */
+  1,			      /* free_if			      */
+  4,			      /* align				      */
+  0,			      /* mic_offset			      */
+  { 0, 0, 0, 0, 0, 0, 0, 0 }, /* is_static, is_static_dstn, has_length,
+				 is_stack_buf, sink_addr, alloc_disp,
+				 is_noncont_src, is_noncont_dst	      */
+  0,			      /* offset				      */
+  0,			      /* size				      */
+  1,			      /* count				      */
+  0,			      /* alloc				      */
+  0,			      /* into				      */
+  0			      /* ptr				      */
+};
+
+/* Pointer to the descriptor of the last loaded shared library.  */
+static void *last_loaded_library = NULL;
+
+/* Pointer and size of the variable, used in __offload_target_host2tgt_p[12]
+   and __offload_target_tgt2host_p[12].  */
+static void *last_var_ptr = NULL;
+static int last_var_size = 0;
+
+
+/* Override the corresponding functions from libgomp.  */
+extern "C" int
+omp_is_initial_device (void) __GOMP_NOTHROW
+{
+  return 0;
+}
+
+extern "C" int32_t
+omp_is_initial_device_ (void)
+{
+  return omp_is_initial_device ();
+}
+
+
+/* Dummy function needed for the initialization of target process during the
+   first call to __offload_offload1.  */
+static void
+__offload_target_init_proc (OFFLOAD ofldt)
+{
+  TRACE ("");
+}
+
+/* Collect addresses of the offload functions and of the global variables from
+   the library descriptor and send them to host.
+   Part 1: Send num_funcs and num_vars to host.  */
+static void
+__offload_target_table_p1 (OFFLOAD ofldt)
+{
+  void ***lib_descr = (void ***) last_loaded_library;
+
+  if (lib_descr == NULL)
+    {
+      TRACE ("");
+      fprintf (stderr, "Error! No shared libraries loaded on target.\n");
+      return;
+    }
+
+  void **func_table_begin = lib_descr[0];
+  void **func_table_end   = lib_descr[1];
+  void **var_table_begin  = lib_descr[2];
+  void **var_table_end    = lib_descr[3];
+
+  /* The func table contains only addresses, the var table contains addresses
+     and corresponding sizes.  */
+  int num_funcs = func_table_end - func_table_begin;
+  int num_vars = (var_table_end - var_table_begin) / 2;
+  TRACE ("(num_funcs = %d, num_vars = %d)", num_funcs, num_vars);
+
+  VarDesc vd1[2] = { vd_tgt2host, vd_tgt2host };
+  vd1[0].ptr = &num_funcs;
+  vd1[0].size = sizeof (num_funcs);
+  vd1[1].ptr = &num_vars;
+  vd1[1].size = sizeof (num_vars);
+  VarDesc2 vd2[2] = { { "num_funcs", 0 }, { "num_vars", 0 } };
+
+  __offload_target_enter (ofldt, 2, vd1, vd2);
+  __offload_target_leave (ofldt);
+}
+
+/* Part 2: Send the table with addresses to host.  */
+static void
+__offload_target_table_p2 (OFFLOAD ofldt)
+{
+  void ***lib_descr = (void ***) last_loaded_library;
+  void **func_table_begin = lib_descr[0];
+  void **func_table_end   = lib_descr[1];
+  void **var_table_begin  = lib_descr[2];
+  void **var_table_end    = lib_descr[3];
+
+  int num_funcs = func_table_end - func_table_begin;
+  int num_vars = (var_table_end - var_table_begin) / 2;
+  int table_size = (num_funcs + 2 * num_vars) * sizeof (void *);
+  void **table = (void **) malloc (table_size);
+  TRACE ("(table_size = %d)", table_size);
+
+  VarDesc vd1;
+  vd1 = vd_tgt2host;
+  vd1.ptr = table;
+  vd1.size = table_size;
+  VarDesc2 vd2 = { "table", 0 };
+
+  __offload_target_enter (ofldt, 1, &vd1, &vd2);
+
+  void **p;
+  int i = 0;
+  for (p = func_table_begin; p < func_table_end; p++, i++)
+    table[i] = *p;
+
+  for (p = var_table_begin; p < var_table_end; p++, i++)
+    table[i] = *p;
+
+  __offload_target_leave (ofldt);
+  free (table);
+}
+
+/* Allocate size bytes and send a pointer to the allocated memory to host.  */
+static void
+__offload_target_alloc (OFFLOAD ofldt)
+{
+  size_t size = 0;
+  void *ptr = NULL;
+
+  VarDesc vd1[2] = { vd_host2tgt, vd_tgt2host };
+  vd1[0].ptr = &size;
+  vd1[0].size = sizeof (size);
+  vd1[1].ptr = &ptr;
+  vd1[1].size = sizeof (void *);
+  VarDesc2 vd2[2] = { { "size", 0 }, { "ptr", 0 } };
+
+  __offload_target_enter (ofldt, 2, vd1, vd2);
+  ptr = malloc (size);
+  TRACE ("(size = %d): ptr = %p", size, ptr);
+  __offload_target_leave (ofldt);
+}
+
+/* Free the memory space pointed to by ptr.  */
+static void
+__offload_target_free (OFFLOAD ofldt)
+{
+  void *ptr = 0;
+
+  VarDesc vd1 = vd_host2tgt;
+  vd1.ptr = &ptr;
+  vd1.size = sizeof (void *);
+  VarDesc2 vd2 = { "ptr", 0 };
+
+  __offload_target_enter (ofldt, 1, &vd1, &vd2);
+  TRACE ("(ptr = %p)", ptr);
+  free (ptr);
+  __offload_target_leave (ofldt);
+}
+
+/* Receive var_size bytes from host and store to var_ptr.
+   Part 1: Receive var_ptr and var_size from host.  */
+static void
+__offload_target_host2tgt_p1 (OFFLOAD ofldt)
+{
+  void *var_ptr = NULL;
+  size_t var_size = 0;
+
+  VarDesc vd1[2] = { vd_host2tgt, vd_host2tgt };
+  vd1[0].ptr = &var_ptr;
+  vd1[0].size = sizeof (void *);
+  vd1[1].ptr = &var_size;
+  vd1[1].size = sizeof (var_size);
+  VarDesc2 vd2[2] = { { "var_ptr", 0 }, { "var_size", 0 } };
+
+  __offload_target_enter (ofldt, 2, vd1, vd2);
+  TRACE ("(var_ptr = %p, var_size = %d)", var_ptr, var_size);
+  last_var_ptr = var_ptr;
+  last_var_size = var_size;
+  __offload_target_leave (ofldt);
+}
+
+/* Part 2: Receive the data from host.  */
+static void
+__offload_target_host2tgt_p2 (OFFLOAD ofldt)
+{
+  TRACE ("(last_var_ptr = %p, last_var_size = %d)",
+	 last_var_ptr, last_var_size);
+
+  VarDesc vd1 = vd_host2tgt;
+  vd1.ptr = last_var_ptr;
+  vd1.size = last_var_size;
+  VarDesc2 vd2 = { "var", 0 };
+
+  __offload_target_enter (ofldt, 1, &vd1, &vd2);
+  __offload_target_leave (ofldt);
+}
+
+/* Send var_size bytes from var_ptr to host.
+   Part 1: Receive var_ptr and var_size from host.  */
+static void
+__offload_target_tgt2host_p1 (OFFLOAD ofldt)
+{
+  void *var_ptr = NULL;
+  size_t var_size = 0;
+
+  VarDesc vd1[2] = { vd_host2tgt, vd_host2tgt };
+  vd1[0].ptr = &var_ptr;
+  vd1[0].size = sizeof (void *);
+  vd1[1].ptr = &var_size;
+  vd1[1].size = sizeof (var_size);
+  VarDesc2 vd2[2] = { { "var_ptr", 0 }, { "var_size", 0 } };
+
+  __offload_target_enter (ofldt, 2, vd1, vd2);
+  TRACE ("(var_ptr = %p, var_size = %d)", var_ptr, var_size);
+  last_var_ptr = var_ptr;
+  last_var_size = var_size;
+  __offload_target_leave (ofldt);
+}
+
+/* Part 2: Send the data to host.  */
+static void
+__offload_target_tgt2host_p2 (OFFLOAD ofldt)
+{
+  TRACE ("(last_var_ptr = %p, last_var_size = %d)",
+	 last_var_ptr, last_var_size);
+
+  VarDesc vd1 = vd_tgt2host;
+  vd1.ptr = last_var_ptr;
+  vd1.size = last_var_size;
+  VarDesc2 vd2 = { "var", 0 };
+
+  __offload_target_enter (ofldt, 1, &vd1, &vd2);
+  __offload_target_leave (ofldt);
+}
+
+/* Call offload function by the address fn_ptr and pass vars_ptr to it.  */
+static void
+__offload_target_run (OFFLOAD ofldt)
+{
+  void *fn_ptr;
+  void *vars_ptr;
+
+  VarDesc vd1[2] = { vd_host2tgt, vd_host2tgt };
+  vd1[0].ptr = &fn_ptr;
+  vd1[0].size = sizeof (void *);
+  vd1[1].ptr = &vars_ptr;
+  vd1[1].size = sizeof (void *);
+  VarDesc2 vd2[2] = { { "fn_ptr", 0 }, { "vars_ptr", 0 } };
+
+  __offload_target_enter (ofldt, 2, vd1, vd2);
+  TRACE ("(fn_ptr = %p, vars_ptr = %p)", fn_ptr, vars_ptr);
+  void (*fn)(void *) = (void (*)(void *)) fn_ptr;
+  fn (vars_ptr);
+  __offload_target_leave (ofldt);
+}
+
+
+/* This should be called from every library with offloading.  */
+extern "C" void
+target_register_lib (const void *target_table)
+{
+  TRACE ("(target_table = %p { %p, %p, %p, %p })", target_table,
+	 ((void **) target_table)[0], ((void **) target_table)[1],
+	 ((void **) target_table)[2], ((void **) target_table)[3]);
+
+  last_loaded_library = (void *) target_table;
+}
+
+/* Use __offload_target_main from liboffload.  */
+int
+main (int argc, char **argv)
+{
+  __offload_target_main ();
+  return 0;
+}
+
+
+/* Register offload_target_main's functions in the liboffload.  */
+
+struct Entry {
+  const char *name;
+  void *func;
+};
+
+#define REGISTER(f)				      \
+extern "C" const Entry __offload_target_##f##_$entry  \
+__attribute__ ((section(".OffloadEntryTable."))) = {  \
+  "__offload_target_"#f,			      \
+  (void *) __offload_target_##f			      \
+}
+REGISTER (init_proc);
+REGISTER (table_p1);
+REGISTER (table_p2);
+REGISTER (alloc);
+REGISTER (free);
+REGISTER (host2tgt_p1);
+REGISTER (host2tgt_p2);
+REGISTER (tgt2host_p1);
+REGISTER (tgt2host_p2);
+REGISTER (run);
+#undef REGISTER
-- 
1.7.1

Re: [PATCH 3/4] Add libgomp plugin for Intel MIC

[Jakub Jelinek]

On Tue, Oct 21, 2014 at 09:24:13PM +0400, Ilya Verbin wrote:
> +/* Add path specified in LD_LIBRARY_PATH to MIC_LD_LIBRARY_PATH, which is
> +   required by liboffloadmic.  */
> +__attribute__((constructor))
> +static void
> +set_mic_lib_path (void)
> +{
> +  const char *ld_lib_path = getenv (LD_LIBRARY_PATH_ENV);
> +  const char *mic_lib_path = getenv (MIC_LD_LIBRARY_PATH_ENV);
> +  char *mic_lib_path_new;
> +
> +  if (!ld_lib_path)
> +    return;
> +
> +  mic_lib_path_new = (char *) malloc ((mic_lib_path ? strlen (mic_lib_path) : 0)
> +				      + strlen (ld_lib_path) + 2);

malloc can fail, SIGSEGV in response to that is not desirable.
Can't you fallback to alloca, or use just alloca, or use alloca
with malloc fallback?
> + 
> +  if (!mic_lib_path)
> +    strcpy (mic_lib_path_new, ld_lib_path);
> +  else
> +    sprintf (mic_lib_path_new, "%s:%s", mic_lib_path, ld_lib_path);
> +  setenv (MIC_LD_LIBRARY_PATH_ENV, mic_lib_path_new, 1);
> +  free (mic_lib_path_new);

> +extern "C" void
> +GOMP_OFFLOAD_register_image (void *host_table, void *target_image)
> +{
> +  TRACE ("(host_table = %p, target_image = %p)", host_table, target_image);
> +
> +  if (num_libraries >= 1000)
> +    {
> +      fprintf (stderr,
> +	       "%s: The number of loaded shared libraries is over 1000!\n",
> +	       __FILE__);
> +      exit (1);

Where does this artificial limit come from?  Using libNNN.so library names?
Can't you use lib%d.so instead?

Also, seeing register_image, shouldn't there be
GOMP_OFFLOAD_unregister_image which would be invoked when the library
containing MIC offloading regions is dlclosed?
One could use __cxa_atexit or similar for that, something that is given
&__dso_handle.  Or is no cleanup necessary?  At least unregistering it
from translation tables, because the same addresses might be reused by a
different shared library?
With dlopen/dlclose in mind, 1000 might be easily reached, consider 10000
times dlopening/dlclosing (perhaps over longer time, by long running daemon)
a shared library containg #pragma omp target region.

> +static int first_init = 1;
> +
> +/* Load offload_target_main on target.  */
> +extern "C" void
> +GOMP_OFFLOAD_init_device (int device)
> +{
> +  TRACE ("");
> +  pthread_mutex_lock (&mutex);
> +  if (first_init)
> +    {
> +      __offload_register_image (&main_target_image);
> +      first_init = 0;
> +    }
> +  pthread_mutex_unlock (&mutex);

pthread_once instead?

	Jakub

Re: [PATCH 3/4] Add libgomp plugin for Intel MIC

[Ilya Verbin]

On 22 Oct 11:22, Jakub Jelinek wrote:
> On Tue, Oct 21, 2014 at 09:24:13PM +0400, Ilya Verbin wrote:
> > +  mic_lib_path_new = (char *) malloc ((mic_lib_path ? strlen (mic_lib_path) : 0)
> > +				      + strlen (ld_lib_path) + 2);
> 
> malloc can fail, SIGSEGV in response to that is not desirable.
> Can't you fallback to alloca, or use just alloca, or use alloca
> with malloc fallback?

I replaced it with alloca.

> > +extern "C" void
> > +GOMP_OFFLOAD_register_image (void *host_table, void *target_image)
> > +{
> > +  TRACE ("(host_table = %p, target_image = %p)", host_table, target_image);
> > +
> > +  if (num_libraries >= 1000)
> > +    {
> > +      fprintf (stderr,
> > +	       "%s: The number of loaded shared libraries is over 1000!\n",
> > +	       __FILE__);
> > +      exit (1);
> 
> Where does this artificial limit come from?  Using libNNN.so library names?
> Can't you use lib%d.so instead?

Yes, it comes from the Image structure (liboffloadmic/runtime/offload_host.h:52)
It must contain a null-terminated name, therefore I need to allocate some space
for the name in plugin's struct TargetImage.  But the structure can't contain
any bytes after the trailing zero and before the actual data.
So, now I extended the name to 10 digits and removed the comparison with 1000.

> Also, seeing register_image, shouldn't there be
> GOMP_OFFLOAD_unregister_image which would be invoked when the library
> containing MIC offloading regions is dlclosed?
> One could use __cxa_atexit or similar for that, something that is given
> &__dso_handle.  Or is no cleanup necessary?  At least unregistering it
> from translation tables, because the same addresses might be reused by a
> different shared library?
> With dlopen/dlclose in mind, 1000 might be easily reached, consider 10000
> times dlopening/dlclosing (perhaps over longer time, by long running daemon)
> a shared library containg #pragma omp target region.

Hmm, previously we've tested only cases when all libraries are loaded before the
first offload.  Offloading from a dlopened library after the call to
gomp_target_init isn't working.  So, this will require some changes in
libgomp/target.c .  Is it ok to fix this bug in a separate patch?
And yes, there should be GOMP_OFFLOAD_unregister_image for this case.

> > +static int first_init = 1;
> > +
> > +/* Load offload_target_main on target.  */
> > +extern "C" void
> > +GOMP_OFFLOAD_init_device (int device)
> > +{
> > +  TRACE ("");
> > +  pthread_mutex_lock (&mutex);
> > +  if (first_init)
> > +    {
> > +      __offload_register_image (&main_target_image);
> > +      first_init = 0;
> > +    }
> > +  pthread_mutex_unlock (&mutex);
> 
> pthread_once instead?

Done.

Patch is updated.

Thanks,
  -- Ilya


---

diff --git a/liboffloadmic/configure.ac b/liboffloadmic/configure.ac
index fb575b3..81fae8f 100644
--- a/liboffloadmic/configure.ac
+++ b/liboffloadmic/configure.ac
@@ -42,6 +42,7 @@ AC_PROG_CC
 AC_PROG_CXX
 AC_CONFIG_FILES([Makefile liboffloadmic_host.spec liboffloadmic_target.spec])
 AM_ENABLE_MULTILIB(, ..)
+AC_CONFIG_SUBDIRS(plugin)
 AC_FUNC_ALLOCA
 AC_CHECK_HEADERS([mm_malloc.h], [], [AC_MSG_ERROR(["Couldn't find mm_malloc.h"])])
 AC_CHECK_FUNCS([__secure_getenv secure_getenv])
diff --git a/liboffloadmic/plugin/Makefile.am b/liboffloadmic/plugin/Makefile.am
new file mode 100644
index 0000000..0baf70d
--- /dev/null
+++ b/liboffloadmic/plugin/Makefile.am
@@ -0,0 +1,123 @@
+# Plugin for offload execution on Intel MIC devices.
+#
+# Copyright (C) 2014 Free Software Foundation, Inc.
+#
+# 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).
+#
+# 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/>.
+
+
+AUTOMAKE_OPTIONS = foreign
+ACLOCAL_AMFLAGS = -I ../.. -I ../../config
+
+# Directories
+build_dir = $(top_builddir)
+source_dir = $(top_srcdir)
+coi_inc_dir = $(top_srcdir)/../include/coi
+myo_inc_dir = $(top_srcdir)/../include/myo
+libgomp_src_dir = $(top_srcdir)/../../libgomp
+libgomp_dir = $(build_dir)/../../libgomp
+liboffload_src_dir = $(top_srcdir)/../runtime
+liboffload_dir = $(top_builddir)/..
+
+# May be used by toolexeclibdir.
+gcc_version := $(shell cat $(top_srcdir)/../../gcc/BASE-VER)
+libsubincludedir = $(libdir)/gcc/$(target_alias)/$(gcc_version)$(MULTISUBDIR)/include
+# Search for main_target_image.h in these directories
+target_prefix_dir = $(libdir)/gcc/$(accel_target)/$(gcc_version)$(MULTISUBDIR)
+target_build_dir = $(accel_search_dir)/$(accel_target)$(MULTISUBDIR)/liboffloadmic/plugin
+target_install_dir = $(accel_search_dir)/lib/gcc/$(accel_target)/$(gcc_version)$(MULTISUBDIR)
+
+if PLUGIN_HOST
+  toolexeclib_LTLIBRARIES = libgomp-plugin-intelmic.la
+  libgomp_plugin_intelmic_la_SOURCES = libgomp-plugin-intelmic.cpp
+  libgomp_plugin_intelmic_la_CPPFLAGS = $(CPPFLAGS) -DLINUX -DCOI_LIBRARY_VERSION=2 -DMYO_SUPPORT -DOFFLOAD_DEBUG=1 -DSEP_SUPPORT -DTIMING_SUPPORT -DHOST_LIBRARY=1 -I$(coi_inc_dir) -I$(myo_inc_dir) -I$(liboffload_src_dir) -I$(libgomp_src_dir) -I$(libgomp_dir) -I$(target_prefix_dir)/include -I$(target_build_dir) -I$(target_install_dir)/include
+  libgomp_plugin_intelmic_la_LDFLAGS = -L$(liboffload_dir)/.libs -loffloadmic_host -version-info 1:0:0
+else # PLUGIN_TARGET
+  plugin_includedir = $(libsubincludedir)
+  plugin_include_HEADERS = main_target_image.h
+  AM_CPPFLAGS = $(CPPFLAGS) -DLINUX -DCOI_LIBRARY_VERSION=2 -DMYO_SUPPORT -DOFFLOAD_DEBUG=1 -DSEP_SUPPORT -DTIMING_SUPPORT -DHOST_LIBRARY=0 -I$(coi_inc_dir) -I$(myo_inc_dir) -I$(liboffload_src_dir) -I$(libgomp_dir)
+  AM_CXXFLAGS = $(CXXFLAGS)
+  AM_LDFLAGS = -L$(liboffload_dir)/.libs -L$(libgomp_dir)/.libs -loffloadmic_target -lcoi_device -lmyo-service -lgomp -rdynamic
+endif
+
+main_target_image.h: offload_target_main
+	@echo -n "const int image_size = " > $@
+	@stat -c '%s' $< >> $@
+	@echo ";" >> $@
+	@echo "struct MainTargetImage {" >> $@
+	@echo "  int64_t size;" >> $@
+	@echo "  char name[sizeof \"offload_target_main\"];" >> $@
+	@echo "  char data[image_size];" >> $@
+	@echo "};" >> $@
+	@echo "extern \"C\" const MainTargetImage main_target_image = {" >> $@
+	@echo "  image_size, \"offload_target_main\"," >> $@
+	@cat $< | xxd -include >> $@
+	@echo "};" >> $@
+
+offload_target_main: $(liboffload_dir)/ofldbegin.o offload_target_main.o $(liboffload_dir)/ofldend.o
+	$(CXX) $(AM_LDFLAGS) $^ -o $@
+
+offload_target_main.o: offload_target_main.cpp
+	$(CXX) $(AM_CXXFLAGS) $(AM_CPPFLAGS) -c $< -o $@
+
+# Work around what appears to be a GNU make bug handling MAKEFLAGS
+# values defined in terms of make variables, as is the case for CC and
+# friends when we are called from the top level Makefile.
+AM_MAKEFLAGS = \
+       "AR_FLAGS=$(AR_FLAGS)" \
+       "CC_FOR_BUILD=$(CC_FOR_BUILD)" \
+       "CFLAGS=$(CFLAGS)" \
+       "CXXFLAGS=$(CXXFLAGS)" \
+       "CFLAGS_FOR_BUILD=$(CFLAGS_FOR_BUILD)" \
+       "CFLAGS_FOR_TARGET=$(CFLAGS_FOR_TARGET)" \
+       "INSTALL=$(INSTALL)" \
+       "INSTALL_DATA=$(INSTALL_DATA)" \
+       "INSTALL_PROGRAM=$(INSTALL_PROGRAM)" \
+       "INSTALL_SCRIPT=$(INSTALL_SCRIPT)" \
+       "JC1FLAGS=$(JC1FLAGS)" \
+       "LDFLAGS=$(LDFLAGS)" \
+       "LIBCFLAGS=$(LIBCFLAGS)" \
+       "LIBCFLAGS_FOR_TARGET=$(LIBCFLAGS_FOR_TARGET)" \
+       "MAKE=$(MAKE)" \
+       "MAKEINFO=$(MAKEINFO) $(MAKEINFOFLAGS)" \
+       "PICFLAG=$(PICFLAG)" \
+       "PICFLAG_FOR_TARGET=$(PICFLAG_FOR_TARGET)" \
+       "SHELL=$(SHELL)" \
+       "RUNTESTFLAGS=$(RUNTESTFLAGS)" \
+       "exec_prefix=$(exec_prefix)" \
+       "infodir=$(infodir)" \
+       "libdir=$(libdir)" \
+       "prefix=$(prefix)" \
+       "includedir=$(includedir)" \
+       "AR=$(AR)" \
+       "AS=$(AS)" \
+       "LD=$(LD)" \
+       "LIBCFLAGS=$(LIBCFLAGS)" \
+       "NM=$(NM)" \
+       "PICFLAG=$(PICFLAG)" \
+       "RANLIB=$(RANLIB)" \
+       "DESTDIR=$(DESTDIR)"
+
+MAKEOVERRIDES =
+
diff --git a/liboffloadmic/plugin/configure.ac b/liboffloadmic/plugin/configure.ac
new file mode 100644
index 0000000..283faad
--- /dev/null
+++ b/liboffloadmic/plugin/configure.ac
@@ -0,0 +1,135 @@
+# Plugin for offload execution on Intel MIC devices.
+#
+# Copyright (C) 2014 Free Software Foundation, Inc.
+#
+# Contributed by Andrey Turetskiy <andrey.turetskiy@intel.com>.
+#
+# This file is part of the GNU OpenMP 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/>.
+
+# Process this file with autoconf to produce a configure script, like so:
+# aclocal -I ../.. -I ../../config && autoconf && automake
+
+AC_PREREQ([2.64])
+AC_INIT([Intel MIC Offload Plugin], [1.0], ,[libgomp-plugin-intelmic])
+
+AC_CONFIG_AUX_DIR(../..)
+
+AC_CANONICAL_SYSTEM
+target_alias=${target_alias-$host_alias}
+AC_SUBST(target_alias)
+
+AM_INIT_AUTOMAKE([1.9.0 foreign no-dist])
+
+AM_MAINTAINER_MODE
+
+AC_PROG_CC
+AC_PROG_CXX
+AC_CONFIG_FILES([Makefile])
+AM_ENABLE_MULTILIB(, ../..)
+
+if test "${multilib}" = "yes"; then
+  multilib_arg="--enable-multilib"
+else
+  multilib_arg=
+fi
+
+# Make sure liboffloadmic is enabled
+case "$enable_liboffloadmic" in
+  host | target)
+    ;;
+  *)
+    AC_MSG_ERROR([Liboffloadmic is disabled]) ;;
+esac
+AM_CONDITIONAL(PLUGIN_HOST, [test x"$enable_liboffloadmic" = xhost])
+
+# Get accel target and path to build or install tree of accel compiler
+accel_search_dir=
+accel_target=
+if test x"$enable_liboffloadmic" = xhost; then
+  for accel in `echo $enable_offload_targets | sed -e 's#,# #g'`; do
+    accel_name=`echo $accel | sed 's/=.*//'`
+    accel_dir=`echo $accel | grep '=' | sed 's/.*=//'`
+    case "$accel_name" in
+      *-intelmic-* | *-intelmicemul-*)
+	accel_target=$accel_name
+	accel_search_dir=$accel_dir
+	;;
+    esac
+  done
+  if test x"$accel_target" = x; then
+    AC_MSG_ERROR([--enable-offload-targets does not contain intelmic target])
+  fi
+fi
+AC_SUBST(accel_search_dir)
+AC_SUBST(accel_target)
+
+AC_MSG_CHECKING([for --enable-version-specific-runtime-libs])
+AC_ARG_ENABLE([version-specific-runtime-libs],
+  AC_HELP_STRING([--enable-version-specific-runtime-libs],
+		 [Specify that runtime libraries should be installed in a compiler-specific directory]),
+  [case "$enableval" in
+    yes) enable_version_specific_runtime_libs=yes ;;
+    no)  enable_version_specific_runtime_libs=no ;;
+    *)   AC_MSG_ERROR([Unknown argument to enable/disable version-specific libs]);;
+   esac],
+  [enable_version_specific_runtime_libs=no])
+AC_MSG_RESULT($enable_version_specific_runtime_libs)
+
+
+# Calculate toolexeclibdir.
+# Also toolexecdir, though it's only used in toolexeclibdir.
+case ${enable_version_specific_runtime_libs} in
+  yes)
+    # Need the gcc compiler version to know where to install libraries
+    # and header files if --enable-version-specific-runtime-libs option
+    # is selected.
+    toolexecdir='$(libdir)/gcc/$(target_alias)'
+    toolexeclibdir='$(toolexecdir)/$(gcc_version)$(MULTISUBDIR)'
+    ;;
+  no)
+    if test -n "$with_cross_host" &&
+       test x"$with_cross_host" != x"no"; then
+      # Install a library built with a cross compiler in tooldir, not libdir.
+      toolexecdir='$(exec_prefix)/$(target_alias)'
+      toolexeclibdir='$(toolexecdir)/lib'
+    else
+      toolexecdir='$(libdir)/gcc-lib/$(target_alias)'
+      toolexeclibdir='$(libdir)'
+    fi
+    multi_os_directory=`$CC -print-multi-os-directory`
+    case $multi_os_directory in
+      .) ;; # Avoid trailing /.
+      *) toolexeclibdir=$toolexeclibdir/$multi_os_directory ;;
+    esac
+    ;;
+esac
+
+AC_LIBTOOL_DLOPEN
+AM_PROG_LIBTOOL
+# Forbid libtool to hardcode RPATH, because we want to be able to specify
+# library search directory using LD_LIBRARY_PATH
+hardcode_into_libs=no
+AC_SUBST(toolexecdir)
+AC_SUBST(toolexeclibdir)
+
+# Must be last
+AC_OUTPUT
diff --git a/liboffloadmic/plugin/libgomp-plugin-intelmic.cpp b/liboffloadmic/plugin/libgomp-plugin-intelmic.cpp
new file mode 100644
index 0000000..9c8b3b4
--- /dev/null
+++ b/liboffloadmic/plugin/libgomp-plugin-intelmic.cpp
@@ -0,0 +1,430 @@
+/* Plugin for offload execution on Intel MIC devices.
+
+   Copyright (C) 2014 Free Software Foundation, Inc.
+
+   Contributed by Ilya Verbin <ilya.verbin@intel.com>.
+
+   This file is part of the GNU OpenMP 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/>.  */
+
+/* Host side part of a libgomp plugin.  */
+
+#include <stdint.h>
+#include <stdio.h>
+#include <stdlib.h>
+#include <string.h>
+#include <utility>
+#include <vector>
+#include <libgomp_target.h>
+#include "compiler_if_host.h"
+#include "main_target_image.h"
+
+#define LD_LIBRARY_PATH_ENV	"LD_LIBRARY_PATH"
+#define MIC_LD_LIBRARY_PATH_ENV	"MIC_LD_LIBRARY_PATH"
+
+#ifdef DEBUG
+#define TRACE(...)					    \
+{							    \
+fprintf (stderr, "HOST:\t%s:%s ", __FILE__, __FUNCTION__);  \
+fprintf (stderr, __VA_ARGS__);				    \
+fprintf (stderr, "\n");					    \
+}
+#else
+#define TRACE { }
+#endif
+
+
+static VarDesc vd_host2tgt = {
+  { 1, 1 },		      /* dst, src			      */
+  { 1, 0 },		      /* in, out			      */
+  1,			      /* alloc_if			      */
+  1,			      /* free_if			      */
+  4,			      /* align				      */
+  0,			      /* mic_offset			      */
+  { 0, 0, 0, 0, 0, 0, 0, 0 }, /* is_static, is_static_dstn, has_length,
+				 is_stack_buf, sink_addr, alloc_disp,
+				 is_noncont_src, is_noncont_dst	      */
+  0,			      /* offset				      */
+  0,			      /* size				      */
+  1,			      /* count				      */
+  0,			      /* alloc				      */
+  0,			      /* into				      */
+  0			      /* ptr				      */
+};
+
+static VarDesc vd_tgt2host = {
+  { 1, 1 },		      /* dst, src			      */
+  { 0, 1 },		      /* in, out			      */
+  1,			      /* alloc_if			      */
+  1,			      /* free_if			      */
+  4,			      /* align				      */
+  0,			      /* mic_offset			      */
+  { 0, 0, 0, 0, 0, 0, 0, 0 }, /* is_static, is_static_dstn, has_length,
+				 is_stack_buf, sink_addr, alloc_disp,
+				 is_noncont_src, is_noncont_dst	      */
+  0,			      /* offset				      */
+  0,			      /* size				      */
+  1,			      /* count				      */
+  0,			      /* alloc				      */
+  0,			      /* into				      */
+  0			      /* ptr				      */
+};
+
+
+/* Total number of shared libraries with offloading to Intel MIC.  */
+static int num_libraries;
+
+/* Pointers to the descriptors, containing pointers to host-side tables and to
+   target images.  */
+static std::vector< std::pair<void *, void *> > lib_descrs;
+
+/* Thread-safe registration of the main image.  */
+static pthread_once_t main_image_is_registered = PTHREAD_ONCE_INIT;
+
+
+/* Add path specified in LD_LIBRARY_PATH to MIC_LD_LIBRARY_PATH, which is
+   required by liboffloadmic.  */
+__attribute__((constructor))
+static void
+set_mic_lib_path (void)
+{
+  const char *ld_lib_path = getenv (LD_LIBRARY_PATH_ENV);
+  const char *mic_lib_path = getenv (MIC_LD_LIBRARY_PATH_ENV);
+  char *mic_lib_path_new;
+
+  if (!ld_lib_path)
+    return;
+
+  mic_lib_path_new = (char *) alloca ((mic_lib_path ? strlen (mic_lib_path) : 0)
+				      + strlen (ld_lib_path) + 2);
+
+  if (!mic_lib_path)
+    strcpy (mic_lib_path_new, ld_lib_path);
+  else
+    sprintf (mic_lib_path_new, "%s:%s", mic_lib_path, ld_lib_path);
+  setenv (MIC_LD_LIBRARY_PATH_ENV, mic_lib_path_new, 1);
+}
+
+extern "C" enum offload_target_type
+GOMP_OFFLOAD_get_type (void)
+{
+  enum offload_target_type res = OFFLOAD_TARGET_TYPE_INTEL_MIC;
+  TRACE ("(): return %d", res);
+  return res;
+}
+
+extern "C" int
+GOMP_OFFLOAD_get_num_devices (void)
+{
+  int res = _Offload_number_of_devices ();
+  TRACE ("(): return %d", res);
+  return res;
+}
+
+/* This should be called from every shared library with offloading.  */
+extern "C" void
+GOMP_OFFLOAD_register_image (void *host_table, void *target_image)
+{
+  TRACE ("(host_table = %p, target_image = %p)", host_table, target_image);
+  lib_descrs.push_back (std::make_pair (host_table, target_image));
+  num_libraries++;
+}
+
+static void
+offload (const char *file, uint64_t line, int device, const char *name,
+	 int num_vars, VarDesc *vars, VarDesc2 *vars2)
+{
+  OFFLOAD ofld = __offload_target_acquire1 (&device, file, line);
+  if (ofld)
+    __offload_offload1 (ofld, name, 0, num_vars, vars, vars2, 0, NULL, NULL);
+  else
+    {
+      fprintf (stderr, "%s:%d: Offload target acquire failed\n", file, line);
+      exit (1);
+    }
+}
+
+static void
+register_main_image ()
+{
+  __offload_register_image (&main_target_image);
+}
+
+/* Load offload_target_main on target.  */
+extern "C" void
+GOMP_OFFLOAD_init_device (int device)
+{
+  TRACE ("");
+  pthread_once (&main_image_is_registered, register_main_image);
+  offload (__FILE__, __LINE__, device, "__offload_target_init_proc", 0,
+	   NULL, NULL);
+}
+
+static void
+get_target_table (int device, int &num_funcs, int &num_vars, void **&table)
+{
+  VarDesc vd1[2] = { vd_tgt2host, vd_tgt2host };
+  vd1[0].ptr = &num_funcs;
+  vd1[0].size = sizeof (num_funcs);
+  vd1[1].ptr = &num_vars;
+  vd1[1].size = sizeof (num_vars);
+  VarDesc2 vd1g[2] = { { "num_funcs", 0 }, { "num_vars", 0 } };
+
+  offload (__FILE__, __LINE__, device, "__offload_target_table_p1", 2,
+	   vd1, vd1g);
+
+  int table_size = num_funcs + 2 * num_vars;
+  if (table_size > 0)
+    {
+      table = new void * [table_size];
+
+      VarDesc vd2;
+      vd2 = vd_tgt2host;
+      vd2.ptr = table;
+      vd2.size = table_size * sizeof (void *);
+      VarDesc2 vd2g = { "table", 0 };
+
+      offload (__FILE__, __LINE__, device, "__offload_target_table_p2", 1,
+	       &vd2, &vd2g);
+    }
+}
+
+static void
+load_lib_and_get_table (int device, int lib_num, mapping_table *&table,
+			int &table_size)
+{
+  struct TargetImage {
+    int64_t size;
+    /* 10 characters is enough for max int value.  */
+    char name[sizeof ("lib0000000000.so")];
+    char data[];
+  } __attribute__ ((packed));
+
+  void ***host_table_descr = (void ***) lib_descrs[lib_num].first;
+  void **host_func_start = host_table_descr[0];
+  void **host_func_end   = host_table_descr[1];
+  void **host_var_start  = host_table_descr[2];
+  void **host_var_end    = host_table_descr[3];
+
+  void **target_image_descr = (void **) lib_descrs[lib_num].second;
+  void *image_start = target_image_descr[0];
+  void *image_end   = target_image_descr[1];
+
+  TRACE ("() host_table_descr { %p, %p, %p, %p }", host_func_start,
+	 host_func_end, host_var_start, host_var_end);
+  TRACE ("() target_image_descr { %p, %p }", image_start, image_end);
+
+  int64_t image_size = (uintptr_t) image_end - (uintptr_t) image_start;
+  TargetImage *image
+    = (TargetImage *) malloc (sizeof (int64_t) + sizeof ("lib0000000000.so")
+			      + image_size);
+  image->size = image_size;
+  sprintf (image->name, "lib%010d.so", lib_num);
+  memcpy (image->data, image_start, image->size);
+
+  TRACE ("() __offload_register_image %s { %p, %d }",
+	 image->name, image_start, image->size);
+  __offload_register_image (image);
+
+  int tgt_num_funcs = 0;
+  int tgt_num_vars = 0;
+  void **tgt_table = NULL;
+  get_target_table (device, tgt_num_funcs, tgt_num_vars, tgt_table);
+  free (image);
+
+  /* The func table contains only addresses, the var table contains addresses
+     and corresponding sizes.  */
+  int host_num_funcs = host_func_end - host_func_start;
+  int host_num_vars  = (host_var_end - host_var_start) / 2;
+  TRACE ("() host_num_funcs = %d, tgt_num_funcs = %d",
+	 host_num_funcs, tgt_num_funcs);
+  TRACE ("() host_num_vars = %d, tgt_num_vars = %d",
+	 host_num_vars, tgt_num_vars);
+  if (host_num_funcs != tgt_num_funcs)
+    {
+      fprintf (stderr, "%s: Can't map target functions\n", __FILE__);
+      exit (1);
+    }
+  if (host_num_vars != tgt_num_vars)
+    {
+      fprintf (stderr, "%s: Can't map target variables\n", __FILE__);
+      exit (1);
+    }
+
+  table = (mapping_table *) realloc (table, (table_size + host_num_funcs
+					     + host_num_vars)
+					    * sizeof (mapping_table));
+  if (table == NULL)
+    {
+      fprintf (stderr, "%s: Can't allocate memory\n", __FILE__);
+      exit (1);
+    }
+
+  for (int i = 0; i < host_num_funcs; i++)
+    {
+      mapping_table t;
+      t.host_start = (uintptr_t) host_func_start[i];
+      t.host_end = t.host_start + 1;
+      t.tgt_start = (uintptr_t) tgt_table[i];
+      t.tgt_end = t.tgt_start + 1;
+
+      TRACE ("() lib %d, func %d:\t0x%llx -- 0x%llx",
+	     lib_num, i, t.host_start, t.tgt_start);
+
+      table[table_size++] = t;
+    }
+
+  for (int i = 0; i < host_num_vars * 2; i += 2)
+    {
+      mapping_table t;
+      t.host_start = (uintptr_t) host_var_start[i];
+      t.host_end = t.host_start + (uintptr_t) host_var_start[i+1];
+      t.tgt_start = (uintptr_t) tgt_table[tgt_num_funcs+i];
+      t.tgt_end = t.tgt_start + (uintptr_t) tgt_table[tgt_num_funcs+i+1];
+
+      TRACE ("() lib %d, var %d:\t0x%llx (%d) -- 0x%llx (%d)", lib_num, i/2,
+	     t.host_start, t.host_end - t.host_start,
+	     t.tgt_start, t.tgt_end - t.tgt_start);
+
+      table[table_size++] = t;
+    }
+
+  delete [] tgt_table;
+}
+
+extern "C" int
+GOMP_OFFLOAD_get_table (int device, void *result)
+{
+  TRACE ("(num_libraries = %d)", num_libraries);
+
+  mapping_table *table = NULL;
+  int table_size = 0;
+
+  for (int i = 0; i < num_libraries; i++)
+    load_lib_and_get_table (device, i, table, table_size);
+
+  *(void **) result = table;
+  return table_size;
+}
+
+extern "C" void *
+GOMP_OFFLOAD_alloc (int device, size_t size)
+{
+  TRACE ("(size = %d)", size);
+
+  void *tgt_ptr;
+  VarDesc vd1[2] = { vd_host2tgt, vd_tgt2host };
+  vd1[0].ptr = &size;
+  vd1[0].size = sizeof (size);
+  vd1[1].ptr = &tgt_ptr;
+  vd1[1].size = sizeof (void *);
+  VarDesc2 vd1g[2] = { { "size", 0 }, { "tgt_ptr", 0 } };
+
+  offload (__FILE__, __LINE__, device, "__offload_target_alloc", 2, vd1, vd1g);
+
+  return tgt_ptr;
+}
+
+extern "C" void
+GOMP_OFFLOAD_free (int device, void *tgt_ptr)
+{
+  TRACE ("(tgt_ptr = %p)", tgt_ptr);
+
+  VarDesc vd1 = vd_host2tgt;
+  vd1.ptr = &tgt_ptr;
+  vd1.size = sizeof (void *);
+  VarDesc2 vd1g = { "tgt_ptr", 0 };
+
+  offload (__FILE__, __LINE__, device, "__offload_target_free", 1, &vd1, &vd1g);
+}
+
+extern "C" void *
+GOMP_OFFLOAD_host2dev (int device, void *tgt_ptr, const void *host_ptr,
+		       size_t size)
+{
+  TRACE ("(tgt_ptr = %p, host_ptr = %p, size = %d)", tgt_ptr, host_ptr, size);
+  if (!size)
+    return tgt_ptr;
+
+  VarDesc vd1[2] = { vd_host2tgt, vd_host2tgt };
+  vd1[0].ptr = &tgt_ptr;
+  vd1[0].size = sizeof (void *);
+  vd1[1].ptr = &size;
+  vd1[1].size = sizeof (size);
+  VarDesc2 vd1g[2] = { { "tgt_ptr", 0 }, { "size", 0 } };
+
+  offload (__FILE__, __LINE__, device, "__offload_target_host2tgt_p1", 2,
+	   vd1, vd1g);
+
+  VarDesc vd2 = vd_host2tgt;
+  vd2.ptr = (void *) host_ptr;
+  vd2.size = size;
+  VarDesc2 vd2g = { "var", 0 };
+
+  offload (__FILE__, __LINE__, device, "__offload_target_host2tgt_p2", 1,
+	   &vd2, &vd2g);
+
+  return tgt_ptr;
+}
+
+extern "C" void *
+GOMP_OFFLOAD_dev2host (int device, void *host_ptr, const void *tgt_ptr,
+		       size_t size)
+{
+  TRACE ("(host_ptr = %p, tgt_ptr = %p, size = %d)", host_ptr, tgt_ptr, size);
+  if (!size)
+    return host_ptr;
+
+  VarDesc vd1[2] = { vd_host2tgt, vd_host2tgt };
+  vd1[0].ptr = &tgt_ptr;
+  vd1[0].size = sizeof (void *);
+  vd1[1].ptr = &size;
+  vd1[1].size = sizeof (size);
+  VarDesc2 vd1g[2] = { { "tgt_ptr", 0 }, { "size", 0 } };
+
+  offload (__FILE__, __LINE__, device, "__offload_target_tgt2host_p1", 2,
+	   vd1, vd1g);
+
+  VarDesc vd2 = vd_tgt2host;
+  vd2.ptr = (void *) host_ptr;
+  vd2.size = size;
+  VarDesc2 vd2g = { "var", 0 };
+
+  offload (__FILE__, __LINE__, device, "__offload_target_tgt2host_p2", 1,
+	   &vd2, &vd2g);
+
+  return host_ptr;
+}
+
+extern "C" void
+GOMP_OFFLOAD_run (int device, void *tgt_fn, void *tgt_vars)
+{
+  TRACE ("(tgt_fn = %p, tgt_vars = %p)", tgt_fn, tgt_vars);
+
+  VarDesc vd1[2] = { vd_host2tgt, vd_host2tgt };
+  vd1[0].ptr = &tgt_fn;
+  vd1[0].size = sizeof (void *);
+  vd1[1].ptr = &tgt_vars;
+  vd1[1].size = sizeof (void *);
+  VarDesc2 vd1g[2] = { { "tgt_fn", 0 }, { "tgt_vars", 0 } };
+
+  offload (__FILE__, __LINE__, device, "__offload_target_run", 2, vd1, vd1g);
+}
diff --git a/liboffloadmic/plugin/offload_target_main.cpp b/liboffloadmic/plugin/offload_target_main.cpp
new file mode 100644
index 0000000..4a2778e
--- /dev/null
+++ b/liboffloadmic/plugin/offload_target_main.cpp
@@ -0,0 +1,366 @@
+/* Plugin for offload execution on Intel MIC devices.
+
+   Copyright (C) 2014 Free Software Foundation, Inc.
+
+   Contributed by Ilya Verbin <ilya.verbin@intel.com>.
+
+   This file is part of the GNU OpenMP 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/>.  */
+
+/* Target side part of a libgomp plugin.  */
+
+#include <stdint.h>
+#include <stdio.h>
+#include <stdlib.h>
+#include "compiler_if_target.h"
+
+
+#ifdef DEBUG
+#define TRACE(...)					      \
+{							      \
+fprintf (stderr, "TARGET:\t%s:%s ", __FILE__, __FUNCTION__);  \
+fprintf (stderr, __VA_ARGS__);				      \
+fprintf (stderr, "\n");					      \
+}
+#else
+#define TRACE { }
+#endif
+
+
+static VarDesc vd_host2tgt = {
+  { 1, 1 },		      /* dst, src			      */
+  { 1, 0 },		      /* in, out			      */
+  1,			      /* alloc_if			      */
+  1,			      /* free_if			      */
+  4,			      /* align				      */
+  0,			      /* mic_offset			      */
+  { 0, 0, 0, 0, 0, 0, 0, 0 }, /* is_static, is_static_dstn, has_length,
+				 is_stack_buf, sink_addr, alloc_disp,
+				 is_noncont_src, is_noncont_dst	      */
+  0,			      /* offset				      */
+  0,			      /* size				      */
+  1,			      /* count				      */
+  0,			      /* alloc				      */
+  0,			      /* into				      */
+  0			      /* ptr				      */
+};
+
+static VarDesc vd_tgt2host = {
+  { 1, 1 },		      /* dst, src			      */
+  { 0, 1 },		      /* in, out			      */
+  1,			      /* alloc_if			      */
+  1,			      /* free_if			      */
+  4,			      /* align				      */
+  0,			      /* mic_offset			      */
+  { 0, 0, 0, 0, 0, 0, 0, 0 }, /* is_static, is_static_dstn, has_length,
+				 is_stack_buf, sink_addr, alloc_disp,
+				 is_noncont_src, is_noncont_dst	      */
+  0,			      /* offset				      */
+  0,			      /* size				      */
+  1,			      /* count				      */
+  0,			      /* alloc				      */
+  0,			      /* into				      */
+  0			      /* ptr				      */
+};
+
+/* Pointer to the descriptor of the last loaded shared library.  */
+static void *last_loaded_library = NULL;
+
+/* Pointer and size of the variable, used in __offload_target_host2tgt_p[12]
+   and __offload_target_tgt2host_p[12].  */
+static void *last_var_ptr = NULL;
+static int last_var_size = 0;
+
+
+/* Override the corresponding functions from libgomp.  */
+extern "C" int
+omp_is_initial_device (void) __GOMP_NOTHROW
+{
+  return 0;
+}
+
+extern "C" int32_t
+omp_is_initial_device_ (void)
+{
+  return omp_is_initial_device ();
+}
+
+
+/* Dummy function needed for the initialization of target process during the
+   first call to __offload_offload1.  */
+static void
+__offload_target_init_proc (OFFLOAD ofldt)
+{
+  TRACE ("");
+}
+
+/* Collect addresses of the offload functions and of the global variables from
+   the library descriptor and send them to host.
+   Part 1: Send num_funcs and num_vars to host.  */
+static void
+__offload_target_table_p1 (OFFLOAD ofldt)
+{
+  void ***lib_descr = (void ***) last_loaded_library;
+
+  if (lib_descr == NULL)
+    {
+      TRACE ("");
+      fprintf (stderr, "Error! No shared libraries loaded on target.\n");
+      return;
+    }
+
+  void **func_table_begin = lib_descr[0];
+  void **func_table_end   = lib_descr[1];
+  void **var_table_begin  = lib_descr[2];
+  void **var_table_end    = lib_descr[3];
+
+  /* The func table contains only addresses, the var table contains addresses
+     and corresponding sizes.  */
+  int num_funcs = func_table_end - func_table_begin;
+  int num_vars = (var_table_end - var_table_begin) / 2;
+  TRACE ("(num_funcs = %d, num_vars = %d)", num_funcs, num_vars);
+
+  VarDesc vd1[2] = { vd_tgt2host, vd_tgt2host };
+  vd1[0].ptr = &num_funcs;
+  vd1[0].size = sizeof (num_funcs);
+  vd1[1].ptr = &num_vars;
+  vd1[1].size = sizeof (num_vars);
+  VarDesc2 vd2[2] = { { "num_funcs", 0 }, { "num_vars", 0 } };
+
+  __offload_target_enter (ofldt, 2, vd1, vd2);
+  __offload_target_leave (ofldt);
+}
+
+/* Part 2: Send the table with addresses to host.  */
+static void
+__offload_target_table_p2 (OFFLOAD ofldt)
+{
+  void ***lib_descr = (void ***) last_loaded_library;
+  void **func_table_begin = lib_descr[0];
+  void **func_table_end   = lib_descr[1];
+  void **var_table_begin  = lib_descr[2];
+  void **var_table_end    = lib_descr[3];
+
+  int num_funcs = func_table_end - func_table_begin;
+  int num_vars = (var_table_end - var_table_begin) / 2;
+  int table_size = (num_funcs + 2 * num_vars) * sizeof (void *);
+  void **table = (void **) malloc (table_size);
+  TRACE ("(table_size = %d)", table_size);
+
+  VarDesc vd1;
+  vd1 = vd_tgt2host;
+  vd1.ptr = table;
+  vd1.size = table_size;
+  VarDesc2 vd2 = { "table", 0 };
+
+  __offload_target_enter (ofldt, 1, &vd1, &vd2);
+
+  void **p;
+  int i = 0;
+  for (p = func_table_begin; p < func_table_end; p++, i++)
+    table[i] = *p;
+
+  for (p = var_table_begin; p < var_table_end; p++, i++)
+    table[i] = *p;
+
+  __offload_target_leave (ofldt);
+  free (table);
+}
+
+/* Allocate size bytes and send a pointer to the allocated memory to host.  */
+static void
+__offload_target_alloc (OFFLOAD ofldt)
+{
+  size_t size = 0;
+  void *ptr = NULL;
+
+  VarDesc vd1[2] = { vd_host2tgt, vd_tgt2host };
+  vd1[0].ptr = &size;
+  vd1[0].size = sizeof (size);
+  vd1[1].ptr = &ptr;
+  vd1[1].size = sizeof (void *);
+  VarDesc2 vd2[2] = { { "size", 0 }, { "ptr", 0 } };
+
+  __offload_target_enter (ofldt, 2, vd1, vd2);
+  ptr = malloc (size);
+  TRACE ("(size = %d): ptr = %p", size, ptr);
+  __offload_target_leave (ofldt);
+}
+
+/* Free the memory space pointed to by ptr.  */
+static void
+__offload_target_free (OFFLOAD ofldt)
+{
+  void *ptr = 0;
+
+  VarDesc vd1 = vd_host2tgt;
+  vd1.ptr = &ptr;
+  vd1.size = sizeof (void *);
+  VarDesc2 vd2 = { "ptr", 0 };
+
+  __offload_target_enter (ofldt, 1, &vd1, &vd2);
+  TRACE ("(ptr = %p)", ptr);
+  free (ptr);
+  __offload_target_leave (ofldt);
+}
+
+/* Receive var_size bytes from host and store to var_ptr.
+   Part 1: Receive var_ptr and var_size from host.  */
+static void
+__offload_target_host2tgt_p1 (OFFLOAD ofldt)
+{
+  void *var_ptr = NULL;
+  size_t var_size = 0;
+
+  VarDesc vd1[2] = { vd_host2tgt, vd_host2tgt };
+  vd1[0].ptr = &var_ptr;
+  vd1[0].size = sizeof (void *);
+  vd1[1].ptr = &var_size;
+  vd1[1].size = sizeof (var_size);
+  VarDesc2 vd2[2] = { { "var_ptr", 0 }, { "var_size", 0 } };
+
+  __offload_target_enter (ofldt, 2, vd1, vd2);
+  TRACE ("(var_ptr = %p, var_size = %d)", var_ptr, var_size);
+  last_var_ptr = var_ptr;
+  last_var_size = var_size;
+  __offload_target_leave (ofldt);
+}
+
+/* Part 2: Receive the data from host.  */
+static void
+__offload_target_host2tgt_p2 (OFFLOAD ofldt)
+{
+  TRACE ("(last_var_ptr = %p, last_var_size = %d)",
+	 last_var_ptr, last_var_size);
+
+  VarDesc vd1 = vd_host2tgt;
+  vd1.ptr = last_var_ptr;
+  vd1.size = last_var_size;
+  VarDesc2 vd2 = { "var", 0 };
+
+  __offload_target_enter (ofldt, 1, &vd1, &vd2);
+  __offload_target_leave (ofldt);
+}
+
+/* Send var_size bytes from var_ptr to host.
+   Part 1: Receive var_ptr and var_size from host.  */
+static void
+__offload_target_tgt2host_p1 (OFFLOAD ofldt)
+{
+  void *var_ptr = NULL;
+  size_t var_size = 0;
+
+  VarDesc vd1[2] = { vd_host2tgt, vd_host2tgt };
+  vd1[0].ptr = &var_ptr;
+  vd1[0].size = sizeof (void *);
+  vd1[1].ptr = &var_size;
+  vd1[1].size = sizeof (var_size);
+  VarDesc2 vd2[2] = { { "var_ptr", 0 }, { "var_size", 0 } };
+
+  __offload_target_enter (ofldt, 2, vd1, vd2);
+  TRACE ("(var_ptr = %p, var_size = %d)", var_ptr, var_size);
+  last_var_ptr = var_ptr;
+  last_var_size = var_size;
+  __offload_target_leave (ofldt);
+}
+
+/* Part 2: Send the data to host.  */
+static void
+__offload_target_tgt2host_p2 (OFFLOAD ofldt)
+{
+  TRACE ("(last_var_ptr = %p, last_var_size = %d)",
+	 last_var_ptr, last_var_size);
+
+  VarDesc vd1 = vd_tgt2host;
+  vd1.ptr = last_var_ptr;
+  vd1.size = last_var_size;
+  VarDesc2 vd2 = { "var", 0 };
+
+  __offload_target_enter (ofldt, 1, &vd1, &vd2);
+  __offload_target_leave (ofldt);
+}
+
+/* Call offload function by the address fn_ptr and pass vars_ptr to it.  */
+static void
+__offload_target_run (OFFLOAD ofldt)
+{
+  void *fn_ptr;
+  void *vars_ptr;
+
+  VarDesc vd1[2] = { vd_host2tgt, vd_host2tgt };
+  vd1[0].ptr = &fn_ptr;
+  vd1[0].size = sizeof (void *);
+  vd1[1].ptr = &vars_ptr;
+  vd1[1].size = sizeof (void *);
+  VarDesc2 vd2[2] = { { "fn_ptr", 0 }, { "vars_ptr", 0 } };
+
+  __offload_target_enter (ofldt, 2, vd1, vd2);
+  TRACE ("(fn_ptr = %p, vars_ptr = %p)", fn_ptr, vars_ptr);
+  void (*fn)(void *) = (void (*)(void *)) fn_ptr;
+  fn (vars_ptr);
+  __offload_target_leave (ofldt);
+}
+
+
+/* This should be called from every library with offloading.  */
+extern "C" void
+target_register_lib (const void *target_table)
+{
+  TRACE ("(target_table = %p { %p, %p, %p, %p })", target_table,
+	 ((void **) target_table)[0], ((void **) target_table)[1],
+	 ((void **) target_table)[2], ((void **) target_table)[3]);
+
+  last_loaded_library = (void *) target_table;
+}
+
+/* Use __offload_target_main from liboffload.  */
+int
+main (int argc, char **argv)
+{
+  __offload_target_main ();
+  return 0;
+}
+
+
+/* Register offload_target_main's functions in the liboffload.  */
+
+struct Entry {
+  const char *name;
+  void *func;
+};
+
+#define REGISTER(f)				      \
+extern "C" const Entry __offload_target_##f##_$entry  \
+__attribute__ ((section(".OffloadEntryTable."))) = {  \
+  "__offload_target_"#f,			      \
+  (void *) __offload_target_##f			      \
+}
+REGISTER (init_proc);
+REGISTER (table_p1);
+REGISTER (table_p2);
+REGISTER (alloc);
+REGISTER (free);
+REGISTER (host2tgt_p1);
+REGISTER (host2tgt_p2);
+REGISTER (tgt2host_p1);
+REGISTER (tgt2host_p2);
+REGISTER (run);
+#undef REGISTER
-- 
1.7.1

Re: [PATCH 3/4] Add libgomp plugin for Intel MIC

[Jakub Jelinek]

On Thu, Oct 23, 2014 at 07:41:12PM +0400, Ilya Verbin wrote:
> > malloc can fail, SIGSEGV in response to that is not desirable.
> > Can't you fallback to alloca, or use just alloca, or use alloca
> > with malloc fallback?
> 
> I replaced it with alloca.

There is a risk if a suid or otherwise priviledge escalated program
uses it and attacker passes huge env vars.
Perhaps use alloca if it is <= 2KB and malloc otherwise, and in that case
if malloc fails, just do a fatal error?

> > Where does this artificial limit come from?  Using libNNN.so library names?
> > Can't you use lib%d.so instead?
> 
> Yes, it comes from the Image structure (liboffloadmic/runtime/offload_host.h:52)
> It must contain a null-terminated name, therefore I need to allocate some space
> for the name in plugin's struct TargetImage.  But the structure can't contain
> any bytes after the trailing zero and before the actual data.
> So, now I extended the name to 10 digits and removed the comparison with 1000.

Ok.

> > Also, seeing register_image, shouldn't there be
> > GOMP_OFFLOAD_unregister_image which would be invoked when the library
> > containing MIC offloading regions is dlclosed?
> > One could use __cxa_atexit or similar for that, something that is given
> > &__dso_handle.  Or is no cleanup necessary?  At least unregistering it
> > from translation tables, because the same addresses might be reused by a
> > different shared library?
> > With dlopen/dlclose in mind, 1000 might be easily reached, consider 10000
> > times dlopening/dlclosing (perhaps over longer time, by long running daemon)
> > a shared library containg #pragma omp target region.
> 
> Hmm, previously we've tested only cases when all libraries are loaded before the
> first offload.  Offloading from a dlopened library after the call to
> gomp_target_init isn't working.  So, this will require some changes in
> libgomp/target.c .  Is it ok to fix this bug in a separate patch?

I guess it can be done incrementally, even during stage3.

	Jakub

Re: [PATCH 3/4] Add libgomp plugin for Intel MIC

[Ilya Verbin]

On 24 Oct 16:35, Jakub Jelinek wrote:
> On Thu, Oct 23, 2014 at 07:41:12PM +0400, Ilya Verbin wrote:
> > > malloc can fail, SIGSEGV in response to that is not desirable.
> > > Can't you fallback to alloca, or use just alloca, or use alloca
> > > with malloc fallback?
> > 
> > I replaced it with alloca.
> 
> There is a risk if a suid or otherwise priviledge escalated program
> uses it and attacker passes huge env vars.
> Perhaps use alloca if it is <= 2KB and malloc otherwise, and in that case
> if malloc fails, just do a fatal error?

Why is this more preferable than just a malloc + fatal error?
This function is executed only once at plugin initialization, therefore no real
performance gain could be achived.

Thanks,
  -- Ilya

Re: [PATCH 3/4] Add libgomp plugin for Intel MIC

[Jakub Jelinek]

On Fri, Oct 24, 2014 at 07:08:44PM +0400, Ilya Verbin wrote:
> On 24 Oct 16:35, Jakub Jelinek wrote:
> > On Thu, Oct 23, 2014 at 07:41:12PM +0400, Ilya Verbin wrote:
> > > > malloc can fail, SIGSEGV in response to that is not desirable.
> > > > Can't you fallback to alloca, or use just alloca, or use alloca
> > > > with malloc fallback?
> > > 
> > > I replaced it with alloca.
> > 
> > There is a risk if a suid or otherwise priviledge escalated program
> > uses it and attacker passes huge env vars.
> > Perhaps use alloca if it is <= 2KB and malloc otherwise, and in that case
> > if malloc fails, just do a fatal error?
> 
> Why is this more preferable than just a malloc + fatal error?
> This function is executed only once at plugin initialization, therefore no real
> performance gain could be achived.

Even if it is executed once, using malloc for short env vars that will be
the 99% of all cases sounds like waste of resources to me.
You already know the strlen of the vars, so it is just a matter of
comparing that and setting a bool flag.

	Jakub

Re: [PATCH 3/4] Add libgomp plugin for Intel MIC

[Ilya Verbin]

On 24 Oct 17:18, Jakub Jelinek wrote:
> On Fri, Oct 24, 2014 at 07:08:44PM +0400, Ilya Verbin wrote:
> > On 24 Oct 16:35, Jakub Jelinek wrote:
> > > On Thu, Oct 23, 2014 at 07:41:12PM +0400, Ilya Verbin wrote:
> > > > > malloc can fail, SIGSEGV in response to that is not desirable.
> > > > > Can't you fallback to alloca, or use just alloca, or use alloca
> > > > > with malloc fallback?
> > > > 
> > > > I replaced it with alloca.
> > > 
> > > There is a risk if a suid or otherwise priviledge escalated program
> > > uses it and attacker passes huge env vars.
> > > Perhaps use alloca if it is <= 2KB and malloc otherwise, and in that case
> > > if malloc fails, just do a fatal error?
> > 
> > Why is this more preferable than just a malloc + fatal error?
> > This function is executed only once at plugin initialization, therefore no real
> > performance gain could be achived.
> 
> Even if it is executed once, using malloc for short env vars that will be
> the 99% of all cases sounds like waste of resources to me.
> You already know the strlen of the vars, so it is just a matter of
> comparing that and setting a bool flag.

Done.  Is it ok?

Thanks,
  -- Ilya


---

diff --git a/liboffloadmic/configure.ac b/liboffloadmic/configure.ac
index fb575b3..81fae8f 100644
--- a/liboffloadmic/configure.ac
+++ b/liboffloadmic/configure.ac
@@ -42,6 +42,7 @@ AC_PROG_CC
 AC_PROG_CXX
 AC_CONFIG_FILES([Makefile liboffloadmic_host.spec liboffloadmic_target.spec])
 AM_ENABLE_MULTILIB(, ..)
+AC_CONFIG_SUBDIRS(plugin)
 AC_FUNC_ALLOCA
 AC_CHECK_HEADERS([mm_malloc.h], [], [AC_MSG_ERROR(["Couldn't find mm_malloc.h"])])
 AC_CHECK_FUNCS([__secure_getenv secure_getenv])
diff --git a/liboffloadmic/plugin/Makefile.am b/liboffloadmic/plugin/Makefile.am
new file mode 100644
index 0000000..0baf70d
--- /dev/null
+++ b/liboffloadmic/plugin/Makefile.am
@@ -0,0 +1,123 @@
+# Plugin for offload execution on Intel MIC devices.
+#
+# Copyright (C) 2014 Free Software Foundation, Inc.
+#
+# 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).
+#
+# 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/>.
+
+
+AUTOMAKE_OPTIONS = foreign
+ACLOCAL_AMFLAGS = -I ../.. -I ../../config
+
+# Directories
+build_dir = $(top_builddir)
+source_dir = $(top_srcdir)
+coi_inc_dir = $(top_srcdir)/../include/coi
+myo_inc_dir = $(top_srcdir)/../include/myo
+libgomp_src_dir = $(top_srcdir)/../../libgomp
+libgomp_dir = $(build_dir)/../../libgomp
+liboffload_src_dir = $(top_srcdir)/../runtime
+liboffload_dir = $(top_builddir)/..
+
+# May be used by toolexeclibdir.
+gcc_version := $(shell cat $(top_srcdir)/../../gcc/BASE-VER)
+libsubincludedir = $(libdir)/gcc/$(target_alias)/$(gcc_version)$(MULTISUBDIR)/include
+# Search for main_target_image.h in these directories
+target_prefix_dir = $(libdir)/gcc/$(accel_target)/$(gcc_version)$(MULTISUBDIR)
+target_build_dir = $(accel_search_dir)/$(accel_target)$(MULTISUBDIR)/liboffloadmic/plugin
+target_install_dir = $(accel_search_dir)/lib/gcc/$(accel_target)/$(gcc_version)$(MULTISUBDIR)
+
+if PLUGIN_HOST
+  toolexeclib_LTLIBRARIES = libgomp-plugin-intelmic.la
+  libgomp_plugin_intelmic_la_SOURCES = libgomp-plugin-intelmic.cpp
+  libgomp_plugin_intelmic_la_CPPFLAGS = $(CPPFLAGS) -DLINUX -DCOI_LIBRARY_VERSION=2 -DMYO_SUPPORT -DOFFLOAD_DEBUG=1 -DSEP_SUPPORT -DTIMING_SUPPORT -DHOST_LIBRARY=1 -I$(coi_inc_dir) -I$(myo_inc_dir) -I$(liboffload_src_dir) -I$(libgomp_src_dir) -I$(libgomp_dir) -I$(target_prefix_dir)/include -I$(target_build_dir) -I$(target_install_dir)/include
+  libgomp_plugin_intelmic_la_LDFLAGS = -L$(liboffload_dir)/.libs -loffloadmic_host -version-info 1:0:0
+else # PLUGIN_TARGET
+  plugin_includedir = $(libsubincludedir)
+  plugin_include_HEADERS = main_target_image.h
+  AM_CPPFLAGS = $(CPPFLAGS) -DLINUX -DCOI_LIBRARY_VERSION=2 -DMYO_SUPPORT -DOFFLOAD_DEBUG=1 -DSEP_SUPPORT -DTIMING_SUPPORT -DHOST_LIBRARY=0 -I$(coi_inc_dir) -I$(myo_inc_dir) -I$(liboffload_src_dir) -I$(libgomp_dir)
+  AM_CXXFLAGS = $(CXXFLAGS)
+  AM_LDFLAGS = -L$(liboffload_dir)/.libs -L$(libgomp_dir)/.libs -loffloadmic_target -lcoi_device -lmyo-service -lgomp -rdynamic
+endif
+
+main_target_image.h: offload_target_main
+	@echo -n "const int image_size = " > $@
+	@stat -c '%s' $< >> $@
+	@echo ";" >> $@
+	@echo "struct MainTargetImage {" >> $@
+	@echo "  int64_t size;" >> $@
+	@echo "  char name[sizeof \"offload_target_main\"];" >> $@
+	@echo "  char data[image_size];" >> $@
+	@echo "};" >> $@
+	@echo "extern \"C\" const MainTargetImage main_target_image = {" >> $@
+	@echo "  image_size, \"offload_target_main\"," >> $@
+	@cat $< | xxd -include >> $@
+	@echo "};" >> $@
+
+offload_target_main: $(liboffload_dir)/ofldbegin.o offload_target_main.o $(liboffload_dir)/ofldend.o
+	$(CXX) $(AM_LDFLAGS) $^ -o $@
+
+offload_target_main.o: offload_target_main.cpp
+	$(CXX) $(AM_CXXFLAGS) $(AM_CPPFLAGS) -c $< -o $@
+
+# Work around what appears to be a GNU make bug handling MAKEFLAGS
+# values defined in terms of make variables, as is the case for CC and
+# friends when we are called from the top level Makefile.
+AM_MAKEFLAGS = \
+       "AR_FLAGS=$(AR_FLAGS)" \
+       "CC_FOR_BUILD=$(CC_FOR_BUILD)" \
+       "CFLAGS=$(CFLAGS)" \
+       "CXXFLAGS=$(CXXFLAGS)" \
+       "CFLAGS_FOR_BUILD=$(CFLAGS_FOR_BUILD)" \
+       "CFLAGS_FOR_TARGET=$(CFLAGS_FOR_TARGET)" \
+       "INSTALL=$(INSTALL)" \
+       "INSTALL_DATA=$(INSTALL_DATA)" \
+       "INSTALL_PROGRAM=$(INSTALL_PROGRAM)" \
+       "INSTALL_SCRIPT=$(INSTALL_SCRIPT)" \
+       "JC1FLAGS=$(JC1FLAGS)" \
+       "LDFLAGS=$(LDFLAGS)" \
+       "LIBCFLAGS=$(LIBCFLAGS)" \
+       "LIBCFLAGS_FOR_TARGET=$(LIBCFLAGS_FOR_TARGET)" \
+       "MAKE=$(MAKE)" \
+       "MAKEINFO=$(MAKEINFO) $(MAKEINFOFLAGS)" \
+       "PICFLAG=$(PICFLAG)" \
+       "PICFLAG_FOR_TARGET=$(PICFLAG_FOR_TARGET)" \
+       "SHELL=$(SHELL)" \
+       "RUNTESTFLAGS=$(RUNTESTFLAGS)" \
+       "exec_prefix=$(exec_prefix)" \
+       "infodir=$(infodir)" \
+       "libdir=$(libdir)" \
+       "prefix=$(prefix)" \
+       "includedir=$(includedir)" \
+       "AR=$(AR)" \
+       "AS=$(AS)" \
+       "LD=$(LD)" \
+       "LIBCFLAGS=$(LIBCFLAGS)" \
+       "NM=$(NM)" \
+       "PICFLAG=$(PICFLAG)" \
+       "RANLIB=$(RANLIB)" \
+       "DESTDIR=$(DESTDIR)"
+
+MAKEOVERRIDES =
+
diff --git a/liboffloadmic/plugin/configure.ac b/liboffloadmic/plugin/configure.ac
new file mode 100644
index 0000000..283faad
--- /dev/null
+++ b/liboffloadmic/plugin/configure.ac
@@ -0,0 +1,135 @@
+# Plugin for offload execution on Intel MIC devices.
+#
+# Copyright (C) 2014 Free Software Foundation, Inc.
+#
+# Contributed by Andrey Turetskiy <andrey.turetskiy@intel.com>.
+#
+# This file is part of the GNU OpenMP 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/>.
+
+# Process this file with autoconf to produce a configure script, like so:
+# aclocal -I ../.. -I ../../config && autoconf && automake
+
+AC_PREREQ([2.64])
+AC_INIT([Intel MIC Offload Plugin], [1.0], ,[libgomp-plugin-intelmic])
+
+AC_CONFIG_AUX_DIR(../..)
+
+AC_CANONICAL_SYSTEM
+target_alias=${target_alias-$host_alias}
+AC_SUBST(target_alias)
+
+AM_INIT_AUTOMAKE([1.9.0 foreign no-dist])
+
+AM_MAINTAINER_MODE
+
+AC_PROG_CC
+AC_PROG_CXX
+AC_CONFIG_FILES([Makefile])
+AM_ENABLE_MULTILIB(, ../..)
+
+if test "${multilib}" = "yes"; then
+  multilib_arg="--enable-multilib"
+else
+  multilib_arg=
+fi
+
+# Make sure liboffloadmic is enabled
+case "$enable_liboffloadmic" in
+  host | target)
+    ;;
+  *)
+    AC_MSG_ERROR([Liboffloadmic is disabled]) ;;
+esac
+AM_CONDITIONAL(PLUGIN_HOST, [test x"$enable_liboffloadmic" = xhost])
+
+# Get accel target and path to build or install tree of accel compiler
+accel_search_dir=
+accel_target=
+if test x"$enable_liboffloadmic" = xhost; then
+  for accel in `echo $enable_offload_targets | sed -e 's#,# #g'`; do
+    accel_name=`echo $accel | sed 's/=.*//'`
+    accel_dir=`echo $accel | grep '=' | sed 's/.*=//'`
+    case "$accel_name" in
+      *-intelmic-* | *-intelmicemul-*)
+	accel_target=$accel_name
+	accel_search_dir=$accel_dir
+	;;
+    esac
+  done
+  if test x"$accel_target" = x; then
+    AC_MSG_ERROR([--enable-offload-targets does not contain intelmic target])
+  fi
+fi
+AC_SUBST(accel_search_dir)
+AC_SUBST(accel_target)
+
+AC_MSG_CHECKING([for --enable-version-specific-runtime-libs])
+AC_ARG_ENABLE([version-specific-runtime-libs],
+  AC_HELP_STRING([--enable-version-specific-runtime-libs],
+		 [Specify that runtime libraries should be installed in a compiler-specific directory]),
+  [case "$enableval" in
+    yes) enable_version_specific_runtime_libs=yes ;;
+    no)  enable_version_specific_runtime_libs=no ;;
+    *)   AC_MSG_ERROR([Unknown argument to enable/disable version-specific libs]);;
+   esac],
+  [enable_version_specific_runtime_libs=no])
+AC_MSG_RESULT($enable_version_specific_runtime_libs)
+
+
+# Calculate toolexeclibdir.
+# Also toolexecdir, though it's only used in toolexeclibdir.
+case ${enable_version_specific_runtime_libs} in
+  yes)
+    # Need the gcc compiler version to know where to install libraries
+    # and header files if --enable-version-specific-runtime-libs option
+    # is selected.
+    toolexecdir='$(libdir)/gcc/$(target_alias)'
+    toolexeclibdir='$(toolexecdir)/$(gcc_version)$(MULTISUBDIR)'
+    ;;
+  no)
+    if test -n "$with_cross_host" &&
+       test x"$with_cross_host" != x"no"; then
+      # Install a library built with a cross compiler in tooldir, not libdir.
+      toolexecdir='$(exec_prefix)/$(target_alias)'
+      toolexeclibdir='$(toolexecdir)/lib'
+    else
+      toolexecdir='$(libdir)/gcc-lib/$(target_alias)'
+      toolexeclibdir='$(libdir)'
+    fi
+    multi_os_directory=`$CC -print-multi-os-directory`
+    case $multi_os_directory in
+      .) ;; # Avoid trailing /.
+      *) toolexeclibdir=$toolexeclibdir/$multi_os_directory ;;
+    esac
+    ;;
+esac
+
+AC_LIBTOOL_DLOPEN
+AM_PROG_LIBTOOL
+# Forbid libtool to hardcode RPATH, because we want to be able to specify
+# library search directory using LD_LIBRARY_PATH
+hardcode_into_libs=no
+AC_SUBST(toolexecdir)
+AC_SUBST(toolexeclibdir)
+
+# Must be last
+AC_OUTPUT
diff --git a/liboffloadmic/plugin/libgomp-plugin-intelmic.cpp b/liboffloadmic/plugin/libgomp-plugin-intelmic.cpp
new file mode 100644
index 0000000..22d8625
--- /dev/null
+++ b/liboffloadmic/plugin/libgomp-plugin-intelmic.cpp
@@ -0,0 +1,448 @@
+/* Plugin for offload execution on Intel MIC devices.
+
+   Copyright (C) 2014 Free Software Foundation, Inc.
+
+   Contributed by Ilya Verbin <ilya.verbin@intel.com>.
+
+   This file is part of the GNU OpenMP 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/>.  */
+
+/* Host side part of a libgomp plugin.  */
+
+#include <stdint.h>
+#include <stdio.h>
+#include <stdlib.h>
+#include <string.h>
+#include <utility>
+#include <vector>
+#include <libgomp_target.h>
+#include "compiler_if_host.h"
+#include "main_target_image.h"
+
+#define LD_LIBRARY_PATH_ENV	"LD_LIBRARY_PATH"
+#define MIC_LD_LIBRARY_PATH_ENV	"MIC_LD_LIBRARY_PATH"
+
+#ifdef DEBUG
+#define TRACE(...)					    \
+{							    \
+fprintf (stderr, "HOST:\t%s:%s ", __FILE__, __FUNCTION__);  \
+fprintf (stderr, __VA_ARGS__);				    \
+fprintf (stderr, "\n");					    \
+}
+#else
+#define TRACE { }
+#endif
+
+
+static VarDesc vd_host2tgt = {
+  { 1, 1 },		      /* dst, src			      */
+  { 1, 0 },		      /* in, out			      */
+  1,			      /* alloc_if			      */
+  1,			      /* free_if			      */
+  4,			      /* align				      */
+  0,			      /* mic_offset			      */
+  { 0, 0, 0, 0, 0, 0, 0, 0 }, /* is_static, is_static_dstn, has_length,
+				 is_stack_buf, sink_addr, alloc_disp,
+				 is_noncont_src, is_noncont_dst	      */
+  0,			      /* offset				      */
+  0,			      /* size				      */
+  1,			      /* count				      */
+  0,			      /* alloc				      */
+  0,			      /* into				      */
+  0			      /* ptr				      */
+};
+
+static VarDesc vd_tgt2host = {
+  { 1, 1 },		      /* dst, src			      */
+  { 0, 1 },		      /* in, out			      */
+  1,			      /* alloc_if			      */
+  1,			      /* free_if			      */
+  4,			      /* align				      */
+  0,			      /* mic_offset			      */
+  { 0, 0, 0, 0, 0, 0, 0, 0 }, /* is_static, is_static_dstn, has_length,
+				 is_stack_buf, sink_addr, alloc_disp,
+				 is_noncont_src, is_noncont_dst	      */
+  0,			      /* offset				      */
+  0,			      /* size				      */
+  1,			      /* count				      */
+  0,			      /* alloc				      */
+  0,			      /* into				      */
+  0			      /* ptr				      */
+};
+
+
+/* Total number of shared libraries with offloading to Intel MIC.  */
+static int num_libraries;
+
+/* Pointers to the descriptors, containing pointers to host-side tables and to
+   target images.  */
+static std::vector< std::pair<void *, void *> > lib_descrs;
+
+/* Thread-safe registration of the main image.  */
+static pthread_once_t main_image_is_registered = PTHREAD_ONCE_INIT;
+
+
+/* Add path specified in LD_LIBRARY_PATH to MIC_LD_LIBRARY_PATH, which is
+   required by liboffloadmic.  */
+__attribute__((constructor))
+static void
+set_mic_lib_path (void)
+{
+  bool use_alloca;
+  const char *ld_lib_path = getenv (LD_LIBRARY_PATH_ENV);
+  const char *mic_lib_path = getenv (MIC_LD_LIBRARY_PATH_ENV);
+  char *mic_lib_path_new;
+  size_t len;
+
+  if (!ld_lib_path)
+    return;
+
+  len = (mic_lib_path ? strlen (mic_lib_path) : 0) + strlen (ld_lib_path) + 2;
+  use_alloca = len <= 2048;
+
+  mic_lib_path_new = (char *) (use_alloca ? alloca (len) : malloc (len));
+  if (!mic_lib_path_new)
+    {
+      fprintf (stderr, "%s: Can't allocate memory\n", __FILE__);
+      exit (1);
+    }
+
+  if (!mic_lib_path)
+    strcpy (mic_lib_path_new, ld_lib_path);
+  else
+    sprintf (mic_lib_path_new, "%s:%s", mic_lib_path, ld_lib_path);
+  setenv (MIC_LD_LIBRARY_PATH_ENV, mic_lib_path_new, 1);
+
+  if (!use_alloca)
+    free (mic_lib_path_new);
+}
+
+extern "C" enum offload_target_type
+GOMP_OFFLOAD_get_type (void)
+{
+  enum offload_target_type res = OFFLOAD_TARGET_TYPE_INTEL_MIC;
+  TRACE ("(): return %d", res);
+  return res;
+}
+
+extern "C" int
+GOMP_OFFLOAD_get_num_devices (void)
+{
+  int res = _Offload_number_of_devices ();
+  TRACE ("(): return %d", res);
+  return res;
+}
+
+/* This should be called from every shared library with offloading.  */
+extern "C" void
+GOMP_OFFLOAD_register_image (void *host_table, void *target_image)
+{
+  TRACE ("(host_table = %p, target_image = %p)", host_table, target_image);
+  lib_descrs.push_back (std::make_pair (host_table, target_image));
+  num_libraries++;
+}
+
+static void
+offload (const char *file, uint64_t line, int device, const char *name,
+	 int num_vars, VarDesc *vars, VarDesc2 *vars2)
+{
+  OFFLOAD ofld = __offload_target_acquire1 (&device, file, line);
+  if (ofld)
+    __offload_offload1 (ofld, name, 0, num_vars, vars, vars2, 0, NULL, NULL);
+  else
+    {
+      fprintf (stderr, "%s:%d: Offload target acquire failed\n", file, line);
+      exit (1);
+    }
+}
+
+static void
+register_main_image ()
+{
+  __offload_register_image (&main_target_image);
+}
+
+/* Load offload_target_main on target.  */
+extern "C" void
+GOMP_OFFLOAD_init_device (int device)
+{
+  TRACE ("");
+  pthread_once (&main_image_is_registered, register_main_image);
+  offload (__FILE__, __LINE__, device, "__offload_target_init_proc", 0,
+	   NULL, NULL);
+}
+
+static void
+get_target_table (int device, int &num_funcs, int &num_vars, void **&table)
+{
+  VarDesc vd1[2] = { vd_tgt2host, vd_tgt2host };
+  vd1[0].ptr = &num_funcs;
+  vd1[0].size = sizeof (num_funcs);
+  vd1[1].ptr = &num_vars;
+  vd1[1].size = sizeof (num_vars);
+  VarDesc2 vd1g[2] = { { "num_funcs", 0 }, { "num_vars", 0 } };
+
+  offload (__FILE__, __LINE__, device, "__offload_target_table_p1", 2,
+	   vd1, vd1g);
+
+  int table_size = num_funcs + 2 * num_vars;
+  if (table_size > 0)
+    {
+      table = new void * [table_size];
+
+      VarDesc vd2;
+      vd2 = vd_tgt2host;
+      vd2.ptr = table;
+      vd2.size = table_size * sizeof (void *);
+      VarDesc2 vd2g = { "table", 0 };
+
+      offload (__FILE__, __LINE__, device, "__offload_target_table_p2", 1,
+	       &vd2, &vd2g);
+    }
+}
+
+static void
+load_lib_and_get_table (int device, int lib_num, mapping_table *&table,
+			int &table_size)
+{
+  struct TargetImage {
+    int64_t size;
+    /* 10 characters is enough for max int value.  */
+    char name[sizeof ("lib0000000000.so")];
+    char data[];
+  } __attribute__ ((packed));
+
+  void ***host_table_descr = (void ***) lib_descrs[lib_num].first;
+  void **host_func_start = host_table_descr[0];
+  void **host_func_end   = host_table_descr[1];
+  void **host_var_start  = host_table_descr[2];
+  void **host_var_end    = host_table_descr[3];
+
+  void **target_image_descr = (void **) lib_descrs[lib_num].second;
+  void *image_start = target_image_descr[0];
+  void *image_end   = target_image_descr[1];
+
+  TRACE ("() host_table_descr { %p, %p, %p, %p }", host_func_start,
+	 host_func_end, host_var_start, host_var_end);
+  TRACE ("() target_image_descr { %p, %p }", image_start, image_end);
+
+  int64_t image_size = (uintptr_t) image_end - (uintptr_t) image_start;
+  TargetImage *image
+    = (TargetImage *) malloc (sizeof (int64_t) + sizeof ("lib0000000000.so")
+			      + image_size);
+  if (!image)
+    {
+      fprintf (stderr, "%s: Can't allocate memory\n", __FILE__);
+      exit (1);
+    }
+
+  image->size = image_size;
+  sprintf (image->name, "lib%010d.so", lib_num);
+  memcpy (image->data, image_start, image->size);
+
+  TRACE ("() __offload_register_image %s { %p, %d }",
+	 image->name, image_start, image->size);
+  __offload_register_image (image);
+
+  int tgt_num_funcs = 0;
+  int tgt_num_vars = 0;
+  void **tgt_table = NULL;
+  get_target_table (device, tgt_num_funcs, tgt_num_vars, tgt_table);
+  free (image);
+
+  /* The func table contains only addresses, the var table contains addresses
+     and corresponding sizes.  */
+  int host_num_funcs = host_func_end - host_func_start;
+  int host_num_vars  = (host_var_end - host_var_start) / 2;
+  TRACE ("() host_num_funcs = %d, tgt_num_funcs = %d",
+	 host_num_funcs, tgt_num_funcs);
+  TRACE ("() host_num_vars = %d, tgt_num_vars = %d",
+	 host_num_vars, tgt_num_vars);
+  if (host_num_funcs != tgt_num_funcs)
+    {
+      fprintf (stderr, "%s: Can't map target functions\n", __FILE__);
+      exit (1);
+    }
+  if (host_num_vars != tgt_num_vars)
+    {
+      fprintf (stderr, "%s: Can't map target variables\n", __FILE__);
+      exit (1);
+    }
+
+  table = (mapping_table *) realloc (table, (table_size + host_num_funcs
+					     + host_num_vars)
+					    * sizeof (mapping_table));
+  if (table == NULL)
+    {
+      fprintf (stderr, "%s: Can't allocate memory\n", __FILE__);
+      exit (1);
+    }
+
+  for (int i = 0; i < host_num_funcs; i++)
+    {
+      mapping_table t;
+      t.host_start = (uintptr_t) host_func_start[i];
+      t.host_end = t.host_start + 1;
+      t.tgt_start = (uintptr_t) tgt_table[i];
+      t.tgt_end = t.tgt_start + 1;
+
+      TRACE ("() lib %d, func %d:\t0x%llx -- 0x%llx",
+	     lib_num, i, t.host_start, t.tgt_start);
+
+      table[table_size++] = t;
+    }
+
+  for (int i = 0; i < host_num_vars * 2; i += 2)
+    {
+      mapping_table t;
+      t.host_start = (uintptr_t) host_var_start[i];
+      t.host_end = t.host_start + (uintptr_t) host_var_start[i+1];
+      t.tgt_start = (uintptr_t) tgt_table[tgt_num_funcs+i];
+      t.tgt_end = t.tgt_start + (uintptr_t) tgt_table[tgt_num_funcs+i+1];
+
+      TRACE ("() lib %d, var %d:\t0x%llx (%d) -- 0x%llx (%d)", lib_num, i/2,
+	     t.host_start, t.host_end - t.host_start,
+	     t.tgt_start, t.tgt_end - t.tgt_start);
+
+      table[table_size++] = t;
+    }
+
+  delete [] tgt_table;
+}
+
+extern "C" int
+GOMP_OFFLOAD_get_table (int device, void *result)
+{
+  TRACE ("(num_libraries = %d)", num_libraries);
+
+  mapping_table *table = NULL;
+  int table_size = 0;
+
+  for (int i = 0; i < num_libraries; i++)
+    load_lib_and_get_table (device, i, table, table_size);
+
+  *(void **) result = table;
+  return table_size;
+}
+
+extern "C" void *
+GOMP_OFFLOAD_alloc (int device, size_t size)
+{
+  TRACE ("(size = %d)", size);
+
+  void *tgt_ptr;
+  VarDesc vd1[2] = { vd_host2tgt, vd_tgt2host };
+  vd1[0].ptr = &size;
+  vd1[0].size = sizeof (size);
+  vd1[1].ptr = &tgt_ptr;
+  vd1[1].size = sizeof (void *);
+  VarDesc2 vd1g[2] = { { "size", 0 }, { "tgt_ptr", 0 } };
+
+  offload (__FILE__, __LINE__, device, "__offload_target_alloc", 2, vd1, vd1g);
+
+  return tgt_ptr;
+}
+
+extern "C" void
+GOMP_OFFLOAD_free (int device, void *tgt_ptr)
+{
+  TRACE ("(tgt_ptr = %p)", tgt_ptr);
+
+  VarDesc vd1 = vd_host2tgt;
+  vd1.ptr = &tgt_ptr;
+  vd1.size = sizeof (void *);
+  VarDesc2 vd1g = { "tgt_ptr", 0 };
+
+  offload (__FILE__, __LINE__, device, "__offload_target_free", 1, &vd1, &vd1g);
+}
+
+extern "C" void *
+GOMP_OFFLOAD_host2dev (int device, void *tgt_ptr, const void *host_ptr,
+		       size_t size)
+{
+  TRACE ("(tgt_ptr = %p, host_ptr = %p, size = %d)", tgt_ptr, host_ptr, size);
+  if (!size)
+    return tgt_ptr;
+
+  VarDesc vd1[2] = { vd_host2tgt, vd_host2tgt };
+  vd1[0].ptr = &tgt_ptr;
+  vd1[0].size = sizeof (void *);
+  vd1[1].ptr = &size;
+  vd1[1].size = sizeof (size);
+  VarDesc2 vd1g[2] = { { "tgt_ptr", 0 }, { "size", 0 } };
+
+  offload (__FILE__, __LINE__, device, "__offload_target_host2tgt_p1", 2,
+	   vd1, vd1g);
+
+  VarDesc vd2 = vd_host2tgt;
+  vd2.ptr = (void *) host_ptr;
+  vd2.size = size;
+  VarDesc2 vd2g = { "var", 0 };
+
+  offload (__FILE__, __LINE__, device, "__offload_target_host2tgt_p2", 1,
+	   &vd2, &vd2g);
+
+  return tgt_ptr;
+}
+
+extern "C" void *
+GOMP_OFFLOAD_dev2host (int device, void *host_ptr, const void *tgt_ptr,
+		       size_t size)
+{
+  TRACE ("(host_ptr = %p, tgt_ptr = %p, size = %d)", host_ptr, tgt_ptr, size);
+  if (!size)
+    return host_ptr;
+
+  VarDesc vd1[2] = { vd_host2tgt, vd_host2tgt };
+  vd1[0].ptr = &tgt_ptr;
+  vd1[0].size = sizeof (void *);
+  vd1[1].ptr = &size;
+  vd1[1].size = sizeof (size);
+  VarDesc2 vd1g[2] = { { "tgt_ptr", 0 }, { "size", 0 } };
+
+  offload (__FILE__, __LINE__, device, "__offload_target_tgt2host_p1", 2,
+	   vd1, vd1g);
+
+  VarDesc vd2 = vd_tgt2host;
+  vd2.ptr = (void *) host_ptr;
+  vd2.size = size;
+  VarDesc2 vd2g = { "var", 0 };
+
+  offload (__FILE__, __LINE__, device, "__offload_target_tgt2host_p2", 1,
+	   &vd2, &vd2g);
+
+  return host_ptr;
+}
+
+extern "C" void
+GOMP_OFFLOAD_run (int device, void *tgt_fn, void *tgt_vars)
+{
+  TRACE ("(tgt_fn = %p, tgt_vars = %p)", tgt_fn, tgt_vars);
+
+  VarDesc vd1[2] = { vd_host2tgt, vd_host2tgt };
+  vd1[0].ptr = &tgt_fn;
+  vd1[0].size = sizeof (void *);
+  vd1[1].ptr = &tgt_vars;
+  vd1[1].size = sizeof (void *);
+  VarDesc2 vd1g[2] = { { "tgt_fn", 0 }, { "tgt_vars", 0 } };
+
+  offload (__FILE__, __LINE__, device, "__offload_target_run", 2, vd1, vd1g);
+}
diff --git a/liboffloadmic/plugin/offload_target_main.cpp b/liboffloadmic/plugin/offload_target_main.cpp
new file mode 100644
index 0000000..4a2778e
--- /dev/null
+++ b/liboffloadmic/plugin/offload_target_main.cpp
@@ -0,0 +1,366 @@
+/* Plugin for offload execution on Intel MIC devices.
+
+   Copyright (C) 2014 Free Software Foundation, Inc.
+
+   Contributed by Ilya Verbin <ilya.verbin@intel.com>.
+
+   This file is part of the GNU OpenMP 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/>.  */
+
+/* Target side part of a libgomp plugin.  */
+
+#include <stdint.h>
+#include <stdio.h>
+#include <stdlib.h>
+#include "compiler_if_target.h"
+
+
+#ifdef DEBUG
+#define TRACE(...)					      \
+{							      \
+fprintf (stderr, "TARGET:\t%s:%s ", __FILE__, __FUNCTION__);  \
+fprintf (stderr, __VA_ARGS__);				      \
+fprintf (stderr, "\n");					      \
+}
+#else
+#define TRACE { }
+#endif
+
+
+static VarDesc vd_host2tgt = {
+  { 1, 1 },		      /* dst, src			      */
+  { 1, 0 },		      /* in, out			      */
+  1,			      /* alloc_if			      */
+  1,			      /* free_if			      */
+  4,			      /* align				      */
+  0,			      /* mic_offset			      */
+  { 0, 0, 0, 0, 0, 0, 0, 0 }, /* is_static, is_static_dstn, has_length,
+				 is_stack_buf, sink_addr, alloc_disp,
+				 is_noncont_src, is_noncont_dst	      */
+  0,			      /* offset				      */
+  0,			      /* size				      */
+  1,			      /* count				      */
+  0,			      /* alloc				      */
+  0,			      /* into				      */
+  0			      /* ptr				      */
+};
+
+static VarDesc vd_tgt2host = {
+  { 1, 1 },		      /* dst, src			      */
+  { 0, 1 },		      /* in, out			      */
+  1,			      /* alloc_if			      */
+  1,			      /* free_if			      */
+  4,			      /* align				      */
+  0,			      /* mic_offset			      */
+  { 0, 0, 0, 0, 0, 0, 0, 0 }, /* is_static, is_static_dstn, has_length,
+				 is_stack_buf, sink_addr, alloc_disp,
+				 is_noncont_src, is_noncont_dst	      */
+  0,			      /* offset				      */
+  0,			      /* size				      */
+  1,			      /* count				      */
+  0,			      /* alloc				      */
+  0,			      /* into				      */
+  0			      /* ptr				      */
+};
+
+/* Pointer to the descriptor of the last loaded shared library.  */
+static void *last_loaded_library = NULL;
+
+/* Pointer and size of the variable, used in __offload_target_host2tgt_p[12]
+   and __offload_target_tgt2host_p[12].  */
+static void *last_var_ptr = NULL;
+static int last_var_size = 0;
+
+
+/* Override the corresponding functions from libgomp.  */
+extern "C" int
+omp_is_initial_device (void) __GOMP_NOTHROW
+{
+  return 0;
+}
+
+extern "C" int32_t
+omp_is_initial_device_ (void)
+{
+  return omp_is_initial_device ();
+}
+
+
+/* Dummy function needed for the initialization of target process during the
+   first call to __offload_offload1.  */
+static void
+__offload_target_init_proc (OFFLOAD ofldt)
+{
+  TRACE ("");
+}
+
+/* Collect addresses of the offload functions and of the global variables from
+   the library descriptor and send them to host.
+   Part 1: Send num_funcs and num_vars to host.  */
+static void
+__offload_target_table_p1 (OFFLOAD ofldt)
+{
+  void ***lib_descr = (void ***) last_loaded_library;
+
+  if (lib_descr == NULL)
+    {
+      TRACE ("");
+      fprintf (stderr, "Error! No shared libraries loaded on target.\n");
+      return;
+    }
+
+  void **func_table_begin = lib_descr[0];
+  void **func_table_end   = lib_descr[1];
+  void **var_table_begin  = lib_descr[2];
+  void **var_table_end    = lib_descr[3];
+
+  /* The func table contains only addresses, the var table contains addresses
+     and corresponding sizes.  */
+  int num_funcs = func_table_end - func_table_begin;
+  int num_vars = (var_table_end - var_table_begin) / 2;
+  TRACE ("(num_funcs = %d, num_vars = %d)", num_funcs, num_vars);
+
+  VarDesc vd1[2] = { vd_tgt2host, vd_tgt2host };
+  vd1[0].ptr = &num_funcs;
+  vd1[0].size = sizeof (num_funcs);
+  vd1[1].ptr = &num_vars;
+  vd1[1].size = sizeof (num_vars);
+  VarDesc2 vd2[2] = { { "num_funcs", 0 }, { "num_vars", 0 } };
+
+  __offload_target_enter (ofldt, 2, vd1, vd2);
+  __offload_target_leave (ofldt);
+}
+
+/* Part 2: Send the table with addresses to host.  */
+static void
+__offload_target_table_p2 (OFFLOAD ofldt)
+{
+  void ***lib_descr = (void ***) last_loaded_library;
+  void **func_table_begin = lib_descr[0];
+  void **func_table_end   = lib_descr[1];
+  void **var_table_begin  = lib_descr[2];
+  void **var_table_end    = lib_descr[3];
+
+  int num_funcs = func_table_end - func_table_begin;
+  int num_vars = (var_table_end - var_table_begin) / 2;
+  int table_size = (num_funcs + 2 * num_vars) * sizeof (void *);
+  void **table = (void **) malloc (table_size);
+  TRACE ("(table_size = %d)", table_size);
+
+  VarDesc vd1;
+  vd1 = vd_tgt2host;
+  vd1.ptr = table;
+  vd1.size = table_size;
+  VarDesc2 vd2 = { "table", 0 };
+
+  __offload_target_enter (ofldt, 1, &vd1, &vd2);
+
+  void **p;
+  int i = 0;
+  for (p = func_table_begin; p < func_table_end; p++, i++)
+    table[i] = *p;
+
+  for (p = var_table_begin; p < var_table_end; p++, i++)
+    table[i] = *p;
+
+  __offload_target_leave (ofldt);
+  free (table);
+}
+
+/* Allocate size bytes and send a pointer to the allocated memory to host.  */
+static void
+__offload_target_alloc (OFFLOAD ofldt)
+{
+  size_t size = 0;
+  void *ptr = NULL;
+
+  VarDesc vd1[2] = { vd_host2tgt, vd_tgt2host };
+  vd1[0].ptr = &size;
+  vd1[0].size = sizeof (size);
+  vd1[1].ptr = &ptr;
+  vd1[1].size = sizeof (void *);
+  VarDesc2 vd2[2] = { { "size", 0 }, { "ptr", 0 } };
+
+  __offload_target_enter (ofldt, 2, vd1, vd2);
+  ptr = malloc (size);
+  TRACE ("(size = %d): ptr = %p", size, ptr);
+  __offload_target_leave (ofldt);
+}
+
+/* Free the memory space pointed to by ptr.  */
+static void
+__offload_target_free (OFFLOAD ofldt)
+{
+  void *ptr = 0;
+
+  VarDesc vd1 = vd_host2tgt;
+  vd1.ptr = &ptr;
+  vd1.size = sizeof (void *);
+  VarDesc2 vd2 = { "ptr", 0 };
+
+  __offload_target_enter (ofldt, 1, &vd1, &vd2);
+  TRACE ("(ptr = %p)", ptr);
+  free (ptr);
+  __offload_target_leave (ofldt);
+}
+
+/* Receive var_size bytes from host and store to var_ptr.
+   Part 1: Receive var_ptr and var_size from host.  */
+static void
+__offload_target_host2tgt_p1 (OFFLOAD ofldt)
+{
+  void *var_ptr = NULL;
+  size_t var_size = 0;
+
+  VarDesc vd1[2] = { vd_host2tgt, vd_host2tgt };
+  vd1[0].ptr = &var_ptr;
+  vd1[0].size = sizeof (void *);
+  vd1[1].ptr = &var_size;
+  vd1[1].size = sizeof (var_size);
+  VarDesc2 vd2[2] = { { "var_ptr", 0 }, { "var_size", 0 } };
+
+  __offload_target_enter (ofldt, 2, vd1, vd2);
+  TRACE ("(var_ptr = %p, var_size = %d)", var_ptr, var_size);
+  last_var_ptr = var_ptr;
+  last_var_size = var_size;
+  __offload_target_leave (ofldt);
+}
+
+/* Part 2: Receive the data from host.  */
+static void
+__offload_target_host2tgt_p2 (OFFLOAD ofldt)
+{
+  TRACE ("(last_var_ptr = %p, last_var_size = %d)",
+	 last_var_ptr, last_var_size);
+
+  VarDesc vd1 = vd_host2tgt;
+  vd1.ptr = last_var_ptr;
+  vd1.size = last_var_size;
+  VarDesc2 vd2 = { "var", 0 };
+
+  __offload_target_enter (ofldt, 1, &vd1, &vd2);
+  __offload_target_leave (ofldt);
+}
+
+/* Send var_size bytes from var_ptr to host.
+   Part 1: Receive var_ptr and var_size from host.  */
+static void
+__offload_target_tgt2host_p1 (OFFLOAD ofldt)
+{
+  void *var_ptr = NULL;
+  size_t var_size = 0;
+
+  VarDesc vd1[2] = { vd_host2tgt, vd_host2tgt };
+  vd1[0].ptr = &var_ptr;
+  vd1[0].size = sizeof (void *);
+  vd1[1].ptr = &var_size;
+  vd1[1].size = sizeof (var_size);
+  VarDesc2 vd2[2] = { { "var_ptr", 0 }, { "var_size", 0 } };
+
+  __offload_target_enter (ofldt, 2, vd1, vd2);
+  TRACE ("(var_ptr = %p, var_size = %d)", var_ptr, var_size);
+  last_var_ptr = var_ptr;
+  last_var_size = var_size;
+  __offload_target_leave (ofldt);
+}
+
+/* Part 2: Send the data to host.  */
+static void
+__offload_target_tgt2host_p2 (OFFLOAD ofldt)
+{
+  TRACE ("(last_var_ptr = %p, last_var_size = %d)",
+	 last_var_ptr, last_var_size);
+
+  VarDesc vd1 = vd_tgt2host;
+  vd1.ptr = last_var_ptr;
+  vd1.size = last_var_size;
+  VarDesc2 vd2 = { "var", 0 };
+
+  __offload_target_enter (ofldt, 1, &vd1, &vd2);
+  __offload_target_leave (ofldt);
+}
+
+/* Call offload function by the address fn_ptr and pass vars_ptr to it.  */
+static void
+__offload_target_run (OFFLOAD ofldt)
+{
+  void *fn_ptr;
+  void *vars_ptr;
+
+  VarDesc vd1[2] = { vd_host2tgt, vd_host2tgt };
+  vd1[0].ptr = &fn_ptr;
+  vd1[0].size = sizeof (void *);
+  vd1[1].ptr = &vars_ptr;
+  vd1[1].size = sizeof (void *);
+  VarDesc2 vd2[2] = { { "fn_ptr", 0 }, { "vars_ptr", 0 } };
+
+  __offload_target_enter (ofldt, 2, vd1, vd2);
+  TRACE ("(fn_ptr = %p, vars_ptr = %p)", fn_ptr, vars_ptr);
+  void (*fn)(void *) = (void (*)(void *)) fn_ptr;
+  fn (vars_ptr);
+  __offload_target_leave (ofldt);
+}
+
+
+/* This should be called from every library with offloading.  */
+extern "C" void
+target_register_lib (const void *target_table)
+{
+  TRACE ("(target_table = %p { %p, %p, %p, %p })", target_table,
+	 ((void **) target_table)[0], ((void **) target_table)[1],
+	 ((void **) target_table)[2], ((void **) target_table)[3]);
+
+  last_loaded_library = (void *) target_table;
+}
+
+/* Use __offload_target_main from liboffload.  */
+int
+main (int argc, char **argv)
+{
+  __offload_target_main ();
+  return 0;
+}
+
+
+/* Register offload_target_main's functions in the liboffload.  */
+
+struct Entry {
+  const char *name;
+  void *func;
+};
+
+#define REGISTER(f)				      \
+extern "C" const Entry __offload_target_##f##_$entry  \
+__attribute__ ((section(".OffloadEntryTable."))) = {  \
+  "__offload_target_"#f,			      \
+  (void *) __offload_target_##f			      \
+}
+REGISTER (init_proc);
+REGISTER (table_p1);
+REGISTER (table_p2);
+REGISTER (alloc);
+REGISTER (free);
+REGISTER (host2tgt_p1);
+REGISTER (host2tgt_p2);
+REGISTER (tgt2host_p1);
+REGISTER (tgt2host_p2);
+REGISTER (run);
+#undef REGISTER
-- 
1.7.1

Re: [PATCH 3/4] Add libgomp plugin for Intel MIC

[Jakub Jelinek]

On Mon, Oct 27, 2014 at 03:15:56PM +0300, Ilya Verbin wrote:

> +  bool use_alloca;
> +  const char *ld_lib_path = getenv (LD_LIBRARY_PATH_ENV);
> +  const char *mic_lib_path = getenv (MIC_LD_LIBRARY_PATH_ENV);
> +  char *mic_lib_path_new;
> +  size_t len;
> +
> +  if (!ld_lib_path)
> +    return;
> +
> +  len = (mic_lib_path ? strlen (mic_lib_path) : 0) + strlen (ld_lib_path) + 2;
> +  use_alloca = len <= 2048;
> +
> +  mic_lib_path_new = (char *) (use_alloca ? alloca (len) : malloc (len));
> +  if (!mic_lib_path_new)
> +    {
> +      fprintf (stderr, "%s: Can't allocate memory\n", __FILE__);
> +      exit (1);
> +    }
> +
> +  if (!mic_lib_path)
> +    strcpy (mic_lib_path_new, ld_lib_path);
> +  else
> +    sprintf (mic_lib_path_new, "%s:%s", mic_lib_path, ld_lib_path);

Oh, one more point, if mic_lib_path is NULL, what is the point
to do the alloca/malloc and string copying?  Can't you just
  setenv (MIC_LD_LIBRARY_PATH_ENV, ld_lib_path, 1);
in that case early?

Otherwise LGTM.

	Jakub

Re: [PATCH 3/4] Add libgomp plugin for Intel MIC

[Ilya Verbin]

On 06 Nov 19:25, Jakub Jelinek wrote:
> Oh, one more point, if mic_lib_path is NULL, what is the point
> to do the alloca/malloc and string copying?  Can't you just
>   setenv (MIC_LD_LIBRARY_PATH_ENV, ld_lib_path, 1);
> in that case early?
> 
> Otherwise LGTM.

Done.

Thanks,
  -- Ilya


---

diff --git a/liboffloadmic/configure.ac b/liboffloadmic/configure.ac
index fb575b3..81fae8f 100644
--- a/liboffloadmic/configure.ac
+++ b/liboffloadmic/configure.ac
@@ -42,6 +42,7 @@ AC_PROG_CC
 AC_PROG_CXX
 AC_CONFIG_FILES([Makefile liboffloadmic_host.spec liboffloadmic_target.spec])
 AM_ENABLE_MULTILIB(, ..)
+AC_CONFIG_SUBDIRS(plugin)
 AC_FUNC_ALLOCA
 AC_CHECK_HEADERS([mm_malloc.h], [], [AC_MSG_ERROR(["Couldn't find mm_malloc.h"])])
 AC_CHECK_FUNCS([__secure_getenv secure_getenv])
diff --git a/liboffloadmic/plugin/Makefile.am b/liboffloadmic/plugin/Makefile.am
new file mode 100644
index 0000000..0baf70d
--- /dev/null
+++ b/liboffloadmic/plugin/Makefile.am
@@ -0,0 +1,123 @@
+# Plugin for offload execution on Intel MIC devices.
+#
+# Copyright (C) 2014 Free Software Foundation, Inc.
+#
+# 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).
+#
+# 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/>.
+
+
+AUTOMAKE_OPTIONS = foreign
+ACLOCAL_AMFLAGS = -I ../.. -I ../../config
+
+# Directories
+build_dir = $(top_builddir)
+source_dir = $(top_srcdir)
+coi_inc_dir = $(top_srcdir)/../include/coi
+myo_inc_dir = $(top_srcdir)/../include/myo
+libgomp_src_dir = $(top_srcdir)/../../libgomp
+libgomp_dir = $(build_dir)/../../libgomp
+liboffload_src_dir = $(top_srcdir)/../runtime
+liboffload_dir = $(top_builddir)/..
+
+# May be used by toolexeclibdir.
+gcc_version := $(shell cat $(top_srcdir)/../../gcc/BASE-VER)
+libsubincludedir = $(libdir)/gcc/$(target_alias)/$(gcc_version)$(MULTISUBDIR)/include
+# Search for main_target_image.h in these directories
+target_prefix_dir = $(libdir)/gcc/$(accel_target)/$(gcc_version)$(MULTISUBDIR)
+target_build_dir = $(accel_search_dir)/$(accel_target)$(MULTISUBDIR)/liboffloadmic/plugin
+target_install_dir = $(accel_search_dir)/lib/gcc/$(accel_target)/$(gcc_version)$(MULTISUBDIR)
+
+if PLUGIN_HOST
+  toolexeclib_LTLIBRARIES = libgomp-plugin-intelmic.la
+  libgomp_plugin_intelmic_la_SOURCES = libgomp-plugin-intelmic.cpp
+  libgomp_plugin_intelmic_la_CPPFLAGS = $(CPPFLAGS) -DLINUX -DCOI_LIBRARY_VERSION=2 -DMYO_SUPPORT -DOFFLOAD_DEBUG=1 -DSEP_SUPPORT -DTIMING_SUPPORT -DHOST_LIBRARY=1 -I$(coi_inc_dir) -I$(myo_inc_dir) -I$(liboffload_src_dir) -I$(libgomp_src_dir) -I$(libgomp_dir) -I$(target_prefix_dir)/include -I$(target_build_dir) -I$(target_install_dir)/include
+  libgomp_plugin_intelmic_la_LDFLAGS = -L$(liboffload_dir)/.libs -loffloadmic_host -version-info 1:0:0
+else # PLUGIN_TARGET
+  plugin_includedir = $(libsubincludedir)
+  plugin_include_HEADERS = main_target_image.h
+  AM_CPPFLAGS = $(CPPFLAGS) -DLINUX -DCOI_LIBRARY_VERSION=2 -DMYO_SUPPORT -DOFFLOAD_DEBUG=1 -DSEP_SUPPORT -DTIMING_SUPPORT -DHOST_LIBRARY=0 -I$(coi_inc_dir) -I$(myo_inc_dir) -I$(liboffload_src_dir) -I$(libgomp_dir)
+  AM_CXXFLAGS = $(CXXFLAGS)
+  AM_LDFLAGS = -L$(liboffload_dir)/.libs -L$(libgomp_dir)/.libs -loffloadmic_target -lcoi_device -lmyo-service -lgomp -rdynamic
+endif
+
+main_target_image.h: offload_target_main
+	@echo -n "const int image_size = " > $@
+	@stat -c '%s' $< >> $@
+	@echo ";" >> $@
+	@echo "struct MainTargetImage {" >> $@
+	@echo "  int64_t size;" >> $@
+	@echo "  char name[sizeof \"offload_target_main\"];" >> $@
+	@echo "  char data[image_size];" >> $@
+	@echo "};" >> $@
+	@echo "extern \"C\" const MainTargetImage main_target_image = {" >> $@
+	@echo "  image_size, \"offload_target_main\"," >> $@
+	@cat $< | xxd -include >> $@
+	@echo "};" >> $@
+
+offload_target_main: $(liboffload_dir)/ofldbegin.o offload_target_main.o $(liboffload_dir)/ofldend.o
+	$(CXX) $(AM_LDFLAGS) $^ -o $@
+
+offload_target_main.o: offload_target_main.cpp
+	$(CXX) $(AM_CXXFLAGS) $(AM_CPPFLAGS) -c $< -o $@
+
+# Work around what appears to be a GNU make bug handling MAKEFLAGS
+# values defined in terms of make variables, as is the case for CC and
+# friends when we are called from the top level Makefile.
+AM_MAKEFLAGS = \
+       "AR_FLAGS=$(AR_FLAGS)" \
+       "CC_FOR_BUILD=$(CC_FOR_BUILD)" \
+       "CFLAGS=$(CFLAGS)" \
+       "CXXFLAGS=$(CXXFLAGS)" \
+       "CFLAGS_FOR_BUILD=$(CFLAGS_FOR_BUILD)" \
+       "CFLAGS_FOR_TARGET=$(CFLAGS_FOR_TARGET)" \
+       "INSTALL=$(INSTALL)" \
+       "INSTALL_DATA=$(INSTALL_DATA)" \
+       "INSTALL_PROGRAM=$(INSTALL_PROGRAM)" \
+       "INSTALL_SCRIPT=$(INSTALL_SCRIPT)" \
+       "JC1FLAGS=$(JC1FLAGS)" \
+       "LDFLAGS=$(LDFLAGS)" \
+       "LIBCFLAGS=$(LIBCFLAGS)" \
+       "LIBCFLAGS_FOR_TARGET=$(LIBCFLAGS_FOR_TARGET)" \
+       "MAKE=$(MAKE)" \
+       "MAKEINFO=$(MAKEINFO) $(MAKEINFOFLAGS)" \
+       "PICFLAG=$(PICFLAG)" \
+       "PICFLAG_FOR_TARGET=$(PICFLAG_FOR_TARGET)" \
+       "SHELL=$(SHELL)" \
+       "RUNTESTFLAGS=$(RUNTESTFLAGS)" \
+       "exec_prefix=$(exec_prefix)" \
+       "infodir=$(infodir)" \
+       "libdir=$(libdir)" \
+       "prefix=$(prefix)" \
+       "includedir=$(includedir)" \
+       "AR=$(AR)" \
+       "AS=$(AS)" \
+       "LD=$(LD)" \
+       "LIBCFLAGS=$(LIBCFLAGS)" \
+       "NM=$(NM)" \
+       "PICFLAG=$(PICFLAG)" \
+       "RANLIB=$(RANLIB)" \
+       "DESTDIR=$(DESTDIR)"
+
+MAKEOVERRIDES =
+
diff --git a/liboffloadmic/plugin/configure.ac b/liboffloadmic/plugin/configure.ac
new file mode 100644
index 0000000..283faad
--- /dev/null
+++ b/liboffloadmic/plugin/configure.ac
@@ -0,0 +1,135 @@
+# Plugin for offload execution on Intel MIC devices.
+#
+# Copyright (C) 2014 Free Software Foundation, Inc.
+#
+# Contributed by Andrey Turetskiy <andrey.turetskiy@intel.com>.
+#
+# This file is part of the GNU OpenMP 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/>.
+
+# Process this file with autoconf to produce a configure script, like so:
+# aclocal -I ../.. -I ../../config && autoconf && automake
+
+AC_PREREQ([2.64])
+AC_INIT([Intel MIC Offload Plugin], [1.0], ,[libgomp-plugin-intelmic])
+
+AC_CONFIG_AUX_DIR(../..)
+
+AC_CANONICAL_SYSTEM
+target_alias=${target_alias-$host_alias}
+AC_SUBST(target_alias)
+
+AM_INIT_AUTOMAKE([1.9.0 foreign no-dist])
+
+AM_MAINTAINER_MODE
+
+AC_PROG_CC
+AC_PROG_CXX
+AC_CONFIG_FILES([Makefile])
+AM_ENABLE_MULTILIB(, ../..)
+
+if test "${multilib}" = "yes"; then
+  multilib_arg="--enable-multilib"
+else
+  multilib_arg=
+fi
+
+# Make sure liboffloadmic is enabled
+case "$enable_liboffloadmic" in
+  host | target)
+    ;;
+  *)
+    AC_MSG_ERROR([Liboffloadmic is disabled]) ;;
+esac
+AM_CONDITIONAL(PLUGIN_HOST, [test x"$enable_liboffloadmic" = xhost])
+
+# Get accel target and path to build or install tree of accel compiler
+accel_search_dir=
+accel_target=
+if test x"$enable_liboffloadmic" = xhost; then
+  for accel in `echo $enable_offload_targets | sed -e 's#,# #g'`; do
+    accel_name=`echo $accel | sed 's/=.*//'`
+    accel_dir=`echo $accel | grep '=' | sed 's/.*=//'`
+    case "$accel_name" in
+      *-intelmic-* | *-intelmicemul-*)
+	accel_target=$accel_name
+	accel_search_dir=$accel_dir
+	;;
+    esac
+  done
+  if test x"$accel_target" = x; then
+    AC_MSG_ERROR([--enable-offload-targets does not contain intelmic target])
+  fi
+fi
+AC_SUBST(accel_search_dir)
+AC_SUBST(accel_target)
+
+AC_MSG_CHECKING([for --enable-version-specific-runtime-libs])
+AC_ARG_ENABLE([version-specific-runtime-libs],
+  AC_HELP_STRING([--enable-version-specific-runtime-libs],
+		 [Specify that runtime libraries should be installed in a compiler-specific directory]),
+  [case "$enableval" in
+    yes) enable_version_specific_runtime_libs=yes ;;
+    no)  enable_version_specific_runtime_libs=no ;;
+    *)   AC_MSG_ERROR([Unknown argument to enable/disable version-specific libs]);;
+   esac],
+  [enable_version_specific_runtime_libs=no])
+AC_MSG_RESULT($enable_version_specific_runtime_libs)
+
+
+# Calculate toolexeclibdir.
+# Also toolexecdir, though it's only used in toolexeclibdir.
+case ${enable_version_specific_runtime_libs} in
+  yes)
+    # Need the gcc compiler version to know where to install libraries
+    # and header files if --enable-version-specific-runtime-libs option
+    # is selected.
+    toolexecdir='$(libdir)/gcc/$(target_alias)'
+    toolexeclibdir='$(toolexecdir)/$(gcc_version)$(MULTISUBDIR)'
+    ;;
+  no)
+    if test -n "$with_cross_host" &&
+       test x"$with_cross_host" != x"no"; then
+      # Install a library built with a cross compiler in tooldir, not libdir.
+      toolexecdir='$(exec_prefix)/$(target_alias)'
+      toolexeclibdir='$(toolexecdir)/lib'
+    else
+      toolexecdir='$(libdir)/gcc-lib/$(target_alias)'
+      toolexeclibdir='$(libdir)'
+    fi
+    multi_os_directory=`$CC -print-multi-os-directory`
+    case $multi_os_directory in
+      .) ;; # Avoid trailing /.
+      *) toolexeclibdir=$toolexeclibdir/$multi_os_directory ;;
+    esac
+    ;;
+esac
+
+AC_LIBTOOL_DLOPEN
+AM_PROG_LIBTOOL
+# Forbid libtool to hardcode RPATH, because we want to be able to specify
+# library search directory using LD_LIBRARY_PATH
+hardcode_into_libs=no
+AC_SUBST(toolexecdir)
+AC_SUBST(toolexeclibdir)
+
+# Must be last
+AC_OUTPUT
diff --git a/liboffloadmic/plugin/libgomp-plugin-intelmic.cpp b/liboffloadmic/plugin/libgomp-plugin-intelmic.cpp
new file mode 100644
index 0000000..28ddbc3
--- /dev/null
+++ b/liboffloadmic/plugin/libgomp-plugin-intelmic.cpp
@@ -0,0 +1,447 @@
+/* Plugin for offload execution on Intel MIC devices.
+
+   Copyright (C) 2014 Free Software Foundation, Inc.
+
+   Contributed by Ilya Verbin <ilya.verbin@intel.com>.
+
+   This file is part of the GNU OpenMP 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/>.  */
+
+/* Host side part of a libgomp plugin.  */
+
+#include <stdint.h>
+#include <stdio.h>
+#include <stdlib.h>
+#include <string.h>
+#include <utility>
+#include <vector>
+#include <libgomp_target.h>
+#include "compiler_if_host.h"
+#include "main_target_image.h"
+
+#define LD_LIBRARY_PATH_ENV	"LD_LIBRARY_PATH"
+#define MIC_LD_LIBRARY_PATH_ENV	"MIC_LD_LIBRARY_PATH"
+
+#ifdef DEBUG
+#define TRACE(...)					    \
+{							    \
+fprintf (stderr, "HOST:\t%s:%s ", __FILE__, __FUNCTION__);  \
+fprintf (stderr, __VA_ARGS__);				    \
+fprintf (stderr, "\n");					    \
+}
+#else
+#define TRACE { }
+#endif
+
+
+static VarDesc vd_host2tgt = {
+  { 1, 1 },		      /* dst, src			      */
+  { 1, 0 },		      /* in, out			      */
+  1,			      /* alloc_if			      */
+  1,			      /* free_if			      */
+  4,			      /* align				      */
+  0,			      /* mic_offset			      */
+  { 0, 0, 0, 0, 0, 0, 0, 0 }, /* is_static, is_static_dstn, has_length,
+				 is_stack_buf, sink_addr, alloc_disp,
+				 is_noncont_src, is_noncont_dst	      */
+  0,			      /* offset				      */
+  0,			      /* size				      */
+  1,			      /* count				      */
+  0,			      /* alloc				      */
+  0,			      /* into				      */
+  0			      /* ptr				      */
+};
+
+static VarDesc vd_tgt2host = {
+  { 1, 1 },		      /* dst, src			      */
+  { 0, 1 },		      /* in, out			      */
+  1,			      /* alloc_if			      */
+  1,			      /* free_if			      */
+  4,			      /* align				      */
+  0,			      /* mic_offset			      */
+  { 0, 0, 0, 0, 0, 0, 0, 0 }, /* is_static, is_static_dstn, has_length,
+				 is_stack_buf, sink_addr, alloc_disp,
+				 is_noncont_src, is_noncont_dst	      */
+  0,			      /* offset				      */
+  0,			      /* size				      */
+  1,			      /* count				      */
+  0,			      /* alloc				      */
+  0,			      /* into				      */
+  0			      /* ptr				      */
+};
+
+
+/* Total number of shared libraries with offloading to Intel MIC.  */
+static int num_libraries;
+
+/* Pointers to the descriptors, containing pointers to host-side tables and to
+   target images.  */
+static std::vector< std::pair<void *, void *> > lib_descrs;
+
+/* Thread-safe registration of the main image.  */
+static pthread_once_t main_image_is_registered = PTHREAD_ONCE_INIT;
+
+
+/* Add path specified in LD_LIBRARY_PATH to MIC_LD_LIBRARY_PATH, which is
+   required by liboffloadmic.  */
+__attribute__((constructor))
+static void
+set_mic_lib_path (void)
+{
+  const char *ld_lib_path = getenv (LD_LIBRARY_PATH_ENV);
+  const char *mic_lib_path = getenv (MIC_LD_LIBRARY_PATH_ENV);
+
+  if (!ld_lib_path)
+    return;
+
+  if (!mic_lib_path)
+    setenv (MIC_LD_LIBRARY_PATH_ENV, ld_lib_path, 1);
+  else
+    {
+      size_t len = strlen (mic_lib_path) + strlen (ld_lib_path) + 2;
+      bool use_alloca = len <= 2048;
+      char *mic_lib_path_new = (char *) (use_alloca ? alloca (len)
+						    : malloc (len));
+      if (!mic_lib_path_new)
+	{
+	  fprintf (stderr, "%s: Can't allocate memory\n", __FILE__);
+	  exit (1);
+	}
+
+      sprintf (mic_lib_path_new, "%s:%s", mic_lib_path, ld_lib_path);
+      setenv (MIC_LD_LIBRARY_PATH_ENV, mic_lib_path_new, 1);
+
+      if (!use_alloca)
+	free (mic_lib_path_new);
+    }
+}
+
+extern "C" enum offload_target_type
+GOMP_OFFLOAD_get_type (void)
+{
+  enum offload_target_type res = OFFLOAD_TARGET_TYPE_INTEL_MIC;
+  TRACE ("(): return %d", res);
+  return res;
+}
+
+extern "C" int
+GOMP_OFFLOAD_get_num_devices (void)
+{
+  int res = _Offload_number_of_devices ();
+  TRACE ("(): return %d", res);
+  return res;
+}
+
+/* This should be called from every shared library with offloading.  */
+extern "C" void
+GOMP_OFFLOAD_register_image (void *host_table, void *target_image)
+{
+  TRACE ("(host_table = %p, target_image = %p)", host_table, target_image);
+  lib_descrs.push_back (std::make_pair (host_table, target_image));
+  num_libraries++;
+}
+
+static void
+offload (const char *file, uint64_t line, int device, const char *name,
+	 int num_vars, VarDesc *vars, VarDesc2 *vars2)
+{
+  OFFLOAD ofld = __offload_target_acquire1 (&device, file, line);
+  if (ofld)
+    __offload_offload1 (ofld, name, 0, num_vars, vars, vars2, 0, NULL, NULL);
+  else
+    {
+      fprintf (stderr, "%s:%d: Offload target acquire failed\n", file, line);
+      exit (1);
+    }
+}
+
+static void
+register_main_image ()
+{
+  __offload_register_image (&main_target_image);
+}
+
+/* Load offload_target_main on target.  */
+extern "C" void
+GOMP_OFFLOAD_init_device (int device)
+{
+  TRACE ("");
+  pthread_once (&main_image_is_registered, register_main_image);
+  offload (__FILE__, __LINE__, device, "__offload_target_init_proc", 0,
+	   NULL, NULL);
+}
+
+static void
+get_target_table (int device, int &num_funcs, int &num_vars, void **&table)
+{
+  VarDesc vd1[2] = { vd_tgt2host, vd_tgt2host };
+  vd1[0].ptr = &num_funcs;
+  vd1[0].size = sizeof (num_funcs);
+  vd1[1].ptr = &num_vars;
+  vd1[1].size = sizeof (num_vars);
+  VarDesc2 vd1g[2] = { { "num_funcs", 0 }, { "num_vars", 0 } };
+
+  offload (__FILE__, __LINE__, device, "__offload_target_table_p1", 2,
+	   vd1, vd1g);
+
+  int table_size = num_funcs + 2 * num_vars;
+  if (table_size > 0)
+    {
+      table = new void * [table_size];
+
+      VarDesc vd2;
+      vd2 = vd_tgt2host;
+      vd2.ptr = table;
+      vd2.size = table_size * sizeof (void *);
+      VarDesc2 vd2g = { "table", 0 };
+
+      offload (__FILE__, __LINE__, device, "__offload_target_table_p2", 1,
+	       &vd2, &vd2g);
+    }
+}
+
+static void
+load_lib_and_get_table (int device, int lib_num, mapping_table *&table,
+			int &table_size)
+{
+  struct TargetImage {
+    int64_t size;
+    /* 10 characters is enough for max int value.  */
+    char name[sizeof ("lib0000000000.so")];
+    char data[];
+  } __attribute__ ((packed));
+
+  void ***host_table_descr = (void ***) lib_descrs[lib_num].first;
+  void **host_func_start = host_table_descr[0];
+  void **host_func_end   = host_table_descr[1];
+  void **host_var_start  = host_table_descr[2];
+  void **host_var_end    = host_table_descr[3];
+
+  void **target_image_descr = (void **) lib_descrs[lib_num].second;
+  void *image_start = target_image_descr[0];
+  void *image_end   = target_image_descr[1];
+
+  TRACE ("() host_table_descr { %p, %p, %p, %p }", host_func_start,
+	 host_func_end, host_var_start, host_var_end);
+  TRACE ("() target_image_descr { %p, %p }", image_start, image_end);
+
+  int64_t image_size = (uintptr_t) image_end - (uintptr_t) image_start;
+  TargetImage *image
+    = (TargetImage *) malloc (sizeof (int64_t) + sizeof ("lib0000000000.so")
+			      + image_size);
+  if (!image)
+    {
+      fprintf (stderr, "%s: Can't allocate memory\n", __FILE__);
+      exit (1);
+    }
+
+  image->size = image_size;
+  sprintf (image->name, "lib%010d.so", lib_num);
+  memcpy (image->data, image_start, image->size);
+
+  TRACE ("() __offload_register_image %s { %p, %d }",
+	 image->name, image_start, image->size);
+  __offload_register_image (image);
+
+  int tgt_num_funcs = 0;
+  int tgt_num_vars = 0;
+  void **tgt_table = NULL;
+  get_target_table (device, tgt_num_funcs, tgt_num_vars, tgt_table);
+  free (image);
+
+  /* The func table contains only addresses, the var table contains addresses
+     and corresponding sizes.  */
+  int host_num_funcs = host_func_end - host_func_start;
+  int host_num_vars  = (host_var_end - host_var_start) / 2;
+  TRACE ("() host_num_funcs = %d, tgt_num_funcs = %d",
+	 host_num_funcs, tgt_num_funcs);
+  TRACE ("() host_num_vars = %d, tgt_num_vars = %d",
+	 host_num_vars, tgt_num_vars);
+  if (host_num_funcs != tgt_num_funcs)
+    {
+      fprintf (stderr, "%s: Can't map target functions\n", __FILE__);
+      exit (1);
+    }
+  if (host_num_vars != tgt_num_vars)
+    {
+      fprintf (stderr, "%s: Can't map target variables\n", __FILE__);
+      exit (1);
+    }
+
+  table = (mapping_table *) realloc (table, (table_size + host_num_funcs
+					     + host_num_vars)
+					    * sizeof (mapping_table));
+  if (table == NULL)
+    {
+      fprintf (stderr, "%s: Can't allocate memory\n", __FILE__);
+      exit (1);
+    }
+
+  for (int i = 0; i < host_num_funcs; i++)
+    {
+      mapping_table t;
+      t.host_start = (uintptr_t) host_func_start[i];
+      t.host_end = t.host_start + 1;
+      t.tgt_start = (uintptr_t) tgt_table[i];
+      t.tgt_end = t.tgt_start + 1;
+
+      TRACE ("() lib %d, func %d:\t0x%llx -- 0x%llx",
+	     lib_num, i, t.host_start, t.tgt_start);
+
+      table[table_size++] = t;
+    }
+
+  for (int i = 0; i < host_num_vars * 2; i += 2)
+    {
+      mapping_table t;
+      t.host_start = (uintptr_t) host_var_start[i];
+      t.host_end = t.host_start + (uintptr_t) host_var_start[i+1];
+      t.tgt_start = (uintptr_t) tgt_table[tgt_num_funcs+i];
+      t.tgt_end = t.tgt_start + (uintptr_t) tgt_table[tgt_num_funcs+i+1];
+
+      TRACE ("() lib %d, var %d:\t0x%llx (%d) -- 0x%llx (%d)", lib_num, i/2,
+	     t.host_start, t.host_end - t.host_start,
+	     t.tgt_start, t.tgt_end - t.tgt_start);
+
+      table[table_size++] = t;
+    }
+
+  delete [] tgt_table;
+}
+
+extern "C" int
+GOMP_OFFLOAD_get_table (int device, void *result)
+{
+  TRACE ("(num_libraries = %d)", num_libraries);
+
+  mapping_table *table = NULL;
+  int table_size = 0;
+
+  for (int i = 0; i < num_libraries; i++)
+    load_lib_and_get_table (device, i, table, table_size);
+
+  *(void **) result = table;
+  return table_size;
+}
+
+extern "C" void *
+GOMP_OFFLOAD_alloc (int device, size_t size)
+{
+  TRACE ("(size = %d)", size);
+
+  void *tgt_ptr;
+  VarDesc vd1[2] = { vd_host2tgt, vd_tgt2host };
+  vd1[0].ptr = &size;
+  vd1[0].size = sizeof (size);
+  vd1[1].ptr = &tgt_ptr;
+  vd1[1].size = sizeof (void *);
+  VarDesc2 vd1g[2] = { { "size", 0 }, { "tgt_ptr", 0 } };
+
+  offload (__FILE__, __LINE__, device, "__offload_target_alloc", 2, vd1, vd1g);
+
+  return tgt_ptr;
+}
+
+extern "C" void
+GOMP_OFFLOAD_free (int device, void *tgt_ptr)
+{
+  TRACE ("(tgt_ptr = %p)", tgt_ptr);
+
+  VarDesc vd1 = vd_host2tgt;
+  vd1.ptr = &tgt_ptr;
+  vd1.size = sizeof (void *);
+  VarDesc2 vd1g = { "tgt_ptr", 0 };
+
+  offload (__FILE__, __LINE__, device, "__offload_target_free", 1, &vd1, &vd1g);
+}
+
+extern "C" void *
+GOMP_OFFLOAD_host2dev (int device, void *tgt_ptr, const void *host_ptr,
+		       size_t size)
+{
+  TRACE ("(tgt_ptr = %p, host_ptr = %p, size = %d)", tgt_ptr, host_ptr, size);
+  if (!size)
+    return tgt_ptr;
+
+  VarDesc vd1[2] = { vd_host2tgt, vd_host2tgt };
+  vd1[0].ptr = &tgt_ptr;
+  vd1[0].size = sizeof (void *);
+  vd1[1].ptr = &size;
+  vd1[1].size = sizeof (size);
+  VarDesc2 vd1g[2] = { { "tgt_ptr", 0 }, { "size", 0 } };
+
+  offload (__FILE__, __LINE__, device, "__offload_target_host2tgt_p1", 2,
+	   vd1, vd1g);
+
+  VarDesc vd2 = vd_host2tgt;
+  vd2.ptr = (void *) host_ptr;
+  vd2.size = size;
+  VarDesc2 vd2g = { "var", 0 };
+
+  offload (__FILE__, __LINE__, device, "__offload_target_host2tgt_p2", 1,
+	   &vd2, &vd2g);
+
+  return tgt_ptr;
+}
+
+extern "C" void *
+GOMP_OFFLOAD_dev2host (int device, void *host_ptr, const void *tgt_ptr,
+		       size_t size)
+{
+  TRACE ("(host_ptr = %p, tgt_ptr = %p, size = %d)", host_ptr, tgt_ptr, size);
+  if (!size)
+    return host_ptr;
+
+  VarDesc vd1[2] = { vd_host2tgt, vd_host2tgt };
+  vd1[0].ptr = &tgt_ptr;
+  vd1[0].size = sizeof (void *);
+  vd1[1].ptr = &size;
+  vd1[1].size = sizeof (size);
+  VarDesc2 vd1g[2] = { { "tgt_ptr", 0 }, { "size", 0 } };
+
+  offload (__FILE__, __LINE__, device, "__offload_target_tgt2host_p1", 2,
+	   vd1, vd1g);
+
+  VarDesc vd2 = vd_tgt2host;
+  vd2.ptr = (void *) host_ptr;
+  vd2.size = size;
+  VarDesc2 vd2g = { "var", 0 };
+
+  offload (__FILE__, __LINE__, device, "__offload_target_tgt2host_p2", 1,
+	   &vd2, &vd2g);
+
+  return host_ptr;
+}
+
+extern "C" void
+GOMP_OFFLOAD_run (int device, void *tgt_fn, void *tgt_vars)
+{
+  TRACE ("(tgt_fn = %p, tgt_vars = %p)", tgt_fn, tgt_vars);
+
+  VarDesc vd1[2] = { vd_host2tgt, vd_host2tgt };
+  vd1[0].ptr = &tgt_fn;
+  vd1[0].size = sizeof (void *);
+  vd1[1].ptr = &tgt_vars;
+  vd1[1].size = sizeof (void *);
+  VarDesc2 vd1g[2] = { { "tgt_fn", 0 }, { "tgt_vars", 0 } };
+
+  offload (__FILE__, __LINE__, device, "__offload_target_run", 2, vd1, vd1g);
+}
diff --git a/liboffloadmic/plugin/offload_target_main.cpp b/liboffloadmic/plugin/offload_target_main.cpp
new file mode 100644
index 0000000..4a2778e
--- /dev/null
+++ b/liboffloadmic/plugin/offload_target_main.cpp
@@ -0,0 +1,366 @@
+/* Plugin for offload execution on Intel MIC devices.
+
+   Copyright (C) 2014 Free Software Foundation, Inc.
+
+   Contributed by Ilya Verbin <ilya.verbin@intel.com>.
+
+   This file is part of the GNU OpenMP 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/>.  */
+
+/* Target side part of a libgomp plugin.  */
+
+#include <stdint.h>
+#include <stdio.h>
+#include <stdlib.h>
+#include "compiler_if_target.h"
+
+
+#ifdef DEBUG
+#define TRACE(...)					      \
+{							      \
+fprintf (stderr, "TARGET:\t%s:%s ", __FILE__, __FUNCTION__);  \
+fprintf (stderr, __VA_ARGS__);				      \
+fprintf (stderr, "\n");					      \
+}
+#else
+#define TRACE { }
+#endif
+
+
+static VarDesc vd_host2tgt = {
+  { 1, 1 },		      /* dst, src			      */
+  { 1, 0 },		      /* in, out			      */
+  1,			      /* alloc_if			      */
+  1,			      /* free_if			      */
+  4,			      /* align				      */
+  0,			      /* mic_offset			      */
+  { 0, 0, 0, 0, 0, 0, 0, 0 }, /* is_static, is_static_dstn, has_length,
+				 is_stack_buf, sink_addr, alloc_disp,
+				 is_noncont_src, is_noncont_dst	      */
+  0,			      /* offset				      */
+  0,			      /* size				      */
+  1,			      /* count				      */
+  0,			      /* alloc				      */
+  0,			      /* into				      */
+  0			      /* ptr				      */
+};
+
+static VarDesc vd_tgt2host = {
+  { 1, 1 },		      /* dst, src			      */
+  { 0, 1 },		      /* in, out			      */
+  1,			      /* alloc_if			      */
+  1,			      /* free_if			      */
+  4,			      /* align				      */
+  0,			      /* mic_offset			      */
+  { 0, 0, 0, 0, 0, 0, 0, 0 }, /* is_static, is_static_dstn, has_length,
+				 is_stack_buf, sink_addr, alloc_disp,
+				 is_noncont_src, is_noncont_dst	      */
+  0,			      /* offset				      */
+  0,			      /* size				      */
+  1,			      /* count				      */
+  0,			      /* alloc				      */
+  0,			      /* into				      */
+  0			      /* ptr				      */
+};
+
+/* Pointer to the descriptor of the last loaded shared library.  */
+static void *last_loaded_library = NULL;
+
+/* Pointer and size of the variable, used in __offload_target_host2tgt_p[12]
+   and __offload_target_tgt2host_p[12].  */
+static void *last_var_ptr = NULL;
+static int last_var_size = 0;
+
+
+/* Override the corresponding functions from libgomp.  */
+extern "C" int
+omp_is_initial_device (void) __GOMP_NOTHROW
+{
+  return 0;
+}
+
+extern "C" int32_t
+omp_is_initial_device_ (void)
+{
+  return omp_is_initial_device ();
+}
+
+
+/* Dummy function needed for the initialization of target process during the
+   first call to __offload_offload1.  */
+static void
+__offload_target_init_proc (OFFLOAD ofldt)
+{
+  TRACE ("");
+}
+
+/* Collect addresses of the offload functions and of the global variables from
+   the library descriptor and send them to host.
+   Part 1: Send num_funcs and num_vars to host.  */
+static void
+__offload_target_table_p1 (OFFLOAD ofldt)
+{
+  void ***lib_descr = (void ***) last_loaded_library;
+
+  if (lib_descr == NULL)
+    {
+      TRACE ("");
+      fprintf (stderr, "Error! No shared libraries loaded on target.\n");
+      return;
+    }
+
+  void **func_table_begin = lib_descr[0];
+  void **func_table_end   = lib_descr[1];
+  void **var_table_begin  = lib_descr[2];
+  void **var_table_end    = lib_descr[3];
+
+  /* The func table contains only addresses, the var table contains addresses
+     and corresponding sizes.  */
+  int num_funcs = func_table_end - func_table_begin;
+  int num_vars = (var_table_end - var_table_begin) / 2;
+  TRACE ("(num_funcs = %d, num_vars = %d)", num_funcs, num_vars);
+
+  VarDesc vd1[2] = { vd_tgt2host, vd_tgt2host };
+  vd1[0].ptr = &num_funcs;
+  vd1[0].size = sizeof (num_funcs);
+  vd1[1].ptr = &num_vars;
+  vd1[1].size = sizeof (num_vars);
+  VarDesc2 vd2[2] = { { "num_funcs", 0 }, { "num_vars", 0 } };
+
+  __offload_target_enter (ofldt, 2, vd1, vd2);
+  __offload_target_leave (ofldt);
+}
+
+/* Part 2: Send the table with addresses to host.  */
+static void
+__offload_target_table_p2 (OFFLOAD ofldt)
+{
+  void ***lib_descr = (void ***) last_loaded_library;
+  void **func_table_begin = lib_descr[0];
+  void **func_table_end   = lib_descr[1];
+  void **var_table_begin  = lib_descr[2];
+  void **var_table_end    = lib_descr[3];
+
+  int num_funcs = func_table_end - func_table_begin;
+  int num_vars = (var_table_end - var_table_begin) / 2;
+  int table_size = (num_funcs + 2 * num_vars) * sizeof (void *);
+  void **table = (void **) malloc (table_size);
+  TRACE ("(table_size = %d)", table_size);
+
+  VarDesc vd1;
+  vd1 = vd_tgt2host;
+  vd1.ptr = table;
+  vd1.size = table_size;
+  VarDesc2 vd2 = { "table", 0 };
+
+  __offload_target_enter (ofldt, 1, &vd1, &vd2);
+
+  void **p;
+  int i = 0;
+  for (p = func_table_begin; p < func_table_end; p++, i++)
+    table[i] = *p;
+
+  for (p = var_table_begin; p < var_table_end; p++, i++)
+    table[i] = *p;
+
+  __offload_target_leave (ofldt);
+  free (table);
+}
+
+/* Allocate size bytes and send a pointer to the allocated memory to host.  */
+static void
+__offload_target_alloc (OFFLOAD ofldt)
+{
+  size_t size = 0;
+  void *ptr = NULL;
+
+  VarDesc vd1[2] = { vd_host2tgt, vd_tgt2host };
+  vd1[0].ptr = &size;
+  vd1[0].size = sizeof (size);
+  vd1[1].ptr = &ptr;
+  vd1[1].size = sizeof (void *);
+  VarDesc2 vd2[2] = { { "size", 0 }, { "ptr", 0 } };
+
+  __offload_target_enter (ofldt, 2, vd1, vd2);
+  ptr = malloc (size);
+  TRACE ("(size = %d): ptr = %p", size, ptr);
+  __offload_target_leave (ofldt);
+}
+
+/* Free the memory space pointed to by ptr.  */
+static void
+__offload_target_free (OFFLOAD ofldt)
+{
+  void *ptr = 0;
+
+  VarDesc vd1 = vd_host2tgt;
+  vd1.ptr = &ptr;
+  vd1.size = sizeof (void *);
+  VarDesc2 vd2 = { "ptr", 0 };
+
+  __offload_target_enter (ofldt, 1, &vd1, &vd2);
+  TRACE ("(ptr = %p)", ptr);
+  free (ptr);
+  __offload_target_leave (ofldt);
+}
+
+/* Receive var_size bytes from host and store to var_ptr.
+   Part 1: Receive var_ptr and var_size from host.  */
+static void
+__offload_target_host2tgt_p1 (OFFLOAD ofldt)
+{
+  void *var_ptr = NULL;
+  size_t var_size = 0;
+
+  VarDesc vd1[2] = { vd_host2tgt, vd_host2tgt };
+  vd1[0].ptr = &var_ptr;
+  vd1[0].size = sizeof (void *);
+  vd1[1].ptr = &var_size;
+  vd1[1].size = sizeof (var_size);
+  VarDesc2 vd2[2] = { { "var_ptr", 0 }, { "var_size", 0 } };
+
+  __offload_target_enter (ofldt, 2, vd1, vd2);
+  TRACE ("(var_ptr = %p, var_size = %d)", var_ptr, var_size);
+  last_var_ptr = var_ptr;
+  last_var_size = var_size;
+  __offload_target_leave (ofldt);
+}
+
+/* Part 2: Receive the data from host.  */
+static void
+__offload_target_host2tgt_p2 (OFFLOAD ofldt)
+{
+  TRACE ("(last_var_ptr = %p, last_var_size = %d)",
+	 last_var_ptr, last_var_size);
+
+  VarDesc vd1 = vd_host2tgt;
+  vd1.ptr = last_var_ptr;
+  vd1.size = last_var_size;
+  VarDesc2 vd2 = { "var", 0 };
+
+  __offload_target_enter (ofldt, 1, &vd1, &vd2);
+  __offload_target_leave (ofldt);
+}
+
+/* Send var_size bytes from var_ptr to host.
+   Part 1: Receive var_ptr and var_size from host.  */
+static void
+__offload_target_tgt2host_p1 (OFFLOAD ofldt)
+{
+  void *var_ptr = NULL;
+  size_t var_size = 0;
+
+  VarDesc vd1[2] = { vd_host2tgt, vd_host2tgt };
+  vd1[0].ptr = &var_ptr;
+  vd1[0].size = sizeof (void *);
+  vd1[1].ptr = &var_size;
+  vd1[1].size = sizeof (var_size);
+  VarDesc2 vd2[2] = { { "var_ptr", 0 }, { "var_size", 0 } };
+
+  __offload_target_enter (ofldt, 2, vd1, vd2);
+  TRACE ("(var_ptr = %p, var_size = %d)", var_ptr, var_size);
+  last_var_ptr = var_ptr;
+  last_var_size = var_size;
+  __offload_target_leave (ofldt);
+}
+
+/* Part 2: Send the data to host.  */
+static void
+__offload_target_tgt2host_p2 (OFFLOAD ofldt)
+{
+  TRACE ("(last_var_ptr = %p, last_var_size = %d)",
+	 last_var_ptr, last_var_size);
+
+  VarDesc vd1 = vd_tgt2host;
+  vd1.ptr = last_var_ptr;
+  vd1.size = last_var_size;
+  VarDesc2 vd2 = { "var", 0 };
+
+  __offload_target_enter (ofldt, 1, &vd1, &vd2);
+  __offload_target_leave (ofldt);
+}
+
+/* Call offload function by the address fn_ptr and pass vars_ptr to it.  */
+static void
+__offload_target_run (OFFLOAD ofldt)
+{
+  void *fn_ptr;
+  void *vars_ptr;
+
+  VarDesc vd1[2] = { vd_host2tgt, vd_host2tgt };
+  vd1[0].ptr = &fn_ptr;
+  vd1[0].size = sizeof (void *);
+  vd1[1].ptr = &vars_ptr;
+  vd1[1].size = sizeof (void *);
+  VarDesc2 vd2[2] = { { "fn_ptr", 0 }, { "vars_ptr", 0 } };
+
+  __offload_target_enter (ofldt, 2, vd1, vd2);
+  TRACE ("(fn_ptr = %p, vars_ptr = %p)", fn_ptr, vars_ptr);
+  void (*fn)(void *) = (void (*)(void *)) fn_ptr;
+  fn (vars_ptr);
+  __offload_target_leave (ofldt);
+}
+
+
+/* This should be called from every library with offloading.  */
+extern "C" void
+target_register_lib (const void *target_table)
+{
+  TRACE ("(target_table = %p { %p, %p, %p, %p })", target_table,
+	 ((void **) target_table)[0], ((void **) target_table)[1],
+	 ((void **) target_table)[2], ((void **) target_table)[3]);
+
+  last_loaded_library = (void *) target_table;
+}
+
+/* Use __offload_target_main from liboffload.  */
+int
+main (int argc, char **argv)
+{
+  __offload_target_main ();
+  return 0;
+}
+
+
+/* Register offload_target_main's functions in the liboffload.  */
+
+struct Entry {
+  const char *name;
+  void *func;
+};
+
+#define REGISTER(f)				      \
+extern "C" const Entry __offload_target_##f##_$entry  \
+__attribute__ ((section(".OffloadEntryTable."))) = {  \
+  "__offload_target_"#f,			      \
+  (void *) __offload_target_##f			      \
+}
+REGISTER (init_proc);
+REGISTER (table_p1);
+REGISTER (table_p2);
+REGISTER (alloc);
+REGISTER (free);
+REGISTER (host2tgt_p1);
+REGISTER (host2tgt_p2);
+REGISTER (tgt2host_p1);
+REGISTER (tgt2host_p2);
+REGISTER (run);
+#undef REGISTER
-- 
1.7.1

Re: [PATCH 3/4] Add libgomp plugin for Intel MIC

[Jakub Jelinek]

On Mon, Nov 10, 2014 at 05:30:38PM +0300, Ilya Verbin wrote:
> On 06 Nov 19:25, Jakub Jelinek wrote:
> > Oh, one more point, if mic_lib_path is NULL, what is the point
> > to do the alloca/malloc and string copying?  Can't you just
> >   setenv (MIC_LD_LIBRARY_PATH_ENV, ld_lib_path, 1);
> > in that case early?
> > 
> > Otherwise LGTM.
> 
> Done.

Ok (with appropriate ChangeLog entry).

	Jakub

Re: [PATCH 3/4] Add libgomp plugin for Intel MIC

[Thomas Schwinge]

[-- Attachment #1: Type: text/plain, Size: 2742 bytes --]

Hi!

On Mon, 10 Nov 2014 17:30:38 +0300, Ilya Verbin <iverbin@gmail.com> wrote:
> --- /dev/null
> +++ b/liboffloadmic/plugin/Makefile.am
> @@ -0,0 +1,123 @@
> +# Plugin for offload execution on Intel MIC devices.

> +libgomp_src_dir = $(top_srcdir)/../../libgomp
> +libgomp_dir = $(build_dir)/../../libgomp

Hmm, I'm not too happy about external (to libgomp) files using (for
example, #include) stuff from libgomp, for the reason given in
<http://news.gmane.org/find-root.php?message_id=%3C87ioishf5z.fsf%40kepler.schwinge.homeip.net%3E>:
it can then easily happen that any such files depend on, for example,
Autoconf definitions which are provided in only one of the instances.
That said, libgomp_target.h as well as omp.h currently are self-contained
(the latter file after having been created from omp.h.in by libgomp's
configure script), so this currently is not an actual problem.


> +  AM_LDFLAGS = -L$(liboffload_dir)/.libs -L$(libgomp_dir)/.libs -loffloadmic_target -lcoi_device -lmyo-service -lgomp -rdynamic

Given that this plugin wishes to link against libgomp, don't we have to
make sure that libgomp has actually been built before that is attempted,
and the following (untested) patch would be required?

diff --git Makefile.def Makefile.def
index 7c8761a..f0a3a91 100644
--- Makefile.def
+++ Makefile.def
@@ -550,7 +550,7 @@ dependencies = { module=configure-target-libvtv; on=all-target-libstdc++-v3; };
 // generated by the libgomp configure.  Unfortunately, due to the use of
 //  recursive make, we can't be that specific.
 dependencies = { module=all-target-libstdc++-v3; on=configure-target-libgomp; };
-dependencies = { module=all-target-liboffloadmic; on=configure-target-libgomp; };
+dependencies = { module=all-target-liboffloadmic; on=all-target-libgomp; };
 
 dependencies = { module=install-target-libgo; on=install-target-libatomic; };
 dependencies = { module=install-target-libgfortran; on=install-target-libquadmath; };
diff --git Makefile.in Makefile.in
index ba5ae4c..8c060b9 100644
--- Makefile.in
+++ Makefile.in
@@ -48884,7 +48884,7 @@ all-stage3-target-libstdc++-v3: maybe-configure-stage3-target-libgomp
 all-stage4-target-libstdc++-v3: maybe-configure-stage4-target-libgomp
 all-stageprofile-target-libstdc++-v3: maybe-configure-stageprofile-target-libgomp
 all-stagefeedback-target-libstdc++-v3: maybe-configure-stagefeedback-target-libgomp
-all-target-liboffloadmic: maybe-configure-target-libgomp
+all-target-liboffloadmic: maybe-all-target-libgomp
 install-target-libgo: maybe-install-target-libatomic
 install-target-libgfortran: maybe-install-target-libquadmath
 install-target-libgfortran: maybe-install-target-libgcc


Grüße,
 Thomas

[-- Attachment #2: signature.asc --]
[-- Type: application/pgp-signature, Size: 472 bytes --]

Re: [PATCH 3/4] Add libgomp plugin for Intel MIC

[Thomas Schwinge]

[-- Attachment #1: Type: text/plain, Size: 4775 bytes --]

Hi!

On Fri, 12 Dec 2014 10:42:30 +0100, I wrote:
> On Mon, 10 Nov 2014 17:30:38 +0300, Ilya Verbin <iverbin@gmail.com> wrote:
> > --- /dev/null
> > +++ b/liboffloadmic/plugin/Makefile.am
> > @@ -0,0 +1,123 @@
> > +# Plugin for offload execution on Intel MIC devices.

> > +  AM_LDFLAGS = -L$(liboffload_dir)/.libs -L$(libgomp_dir)/.libs -loffloadmic_target -lcoi_device -lmyo-service -lgomp -rdynamic
> 
> Given that this plugin wishes to link against libgomp, don't we have to
> make sure that libgomp has actually been built before that is attempted,
> and the following (untested) patch would be required?

As shown by a »make -j1« build: yes, we have to.  As obvious, committed
to trunk in r219344:

commit fb5eef67f6b041cd0bc4f1f8d62c1a000d59f497
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date:   Thu Jan 8 14:46:22 2015 +0000

    liboffloadmic/plugin: Depend on libgomp being built.
    
        [...]
        Making all in plugin
        make[6]: Entering directory `[...]/build-gcc-offload-x86_64-intelmicemul-linux-gnu/x86_64-intelmicemul-linux-gnu/32/liboffloadmic/plugin'
        [...]
        [...]/build-gcc-offload-x86_64-intelmicemul-linux-gnu/./gcc/xg++ -B[...]/build-gcc-offload-x86_64-intelmicemul-linux-gnu/./gcc/ -nostdinc++ -nostdinc++ -I[...]/build-gcc-offload-x86_64-intelmicemul-linux-gnu/x86_64-intelmicemul-linux-gnu/32/libstdc++-v3/include/x86_64-intelmicemul-linux-gnu -I[...]/build-gcc-offload-x86_64-intelmicemul-linux-gnu/x86_64-intelmicemul-linux-gnu/32/libstdc++-v3/include -I[...]/source-gcc/libstdc++-v3/libsupc++ -I[...]/source-gcc/libstdc++-v3/include/backward -I[...]/source-gcc/libstdc++-v3/testsuite/util -L[...]/build-gcc-offload-x86_64-intelmicemul-linux-gnu/x86_64-intelmicemul-linux-gnu/32/libstdc++-v3/src -L[...]/build-gcc-offload-x86_64-intelmicemul-linux-gnu/x86_64-intelmicemul-linux-gnu/32/libstdc++-v3/src/.libs -L[...]/build-gcc-offload-x86_64-intelmicemul-linux-gnu/x86_64-intelmicemul-linux-gnu/32/libstdc++-v3/libsupc++/.libs -B/x86_64-intelmicemul-linux-gnu/bin/ -B/x86_64-intelmicemul-linux-gnu/lib/ -isystem /x86_64-intelmicemul-linux-gnu/include -isystem /x86_64-intelmicemul-linux-gnu/sys-include  -m32 -L./../.libs -L./../../libgomp/.libs -loffloadmic_target -lcoi_device -lmyo-service -lgomp -rdynamic ../ofldbegin.o offload_target_main.o ../ofldend.o -o offload_target_main
        /usr/bin/ld: cannot find -lgomp
        collect2: error: ld returned 1 exit status
    
    	* Makefile.def (dependencies) <all-target-liboffloadmic>: Depend on
    	all-target-libgomp.
    	* Makefile.in: Regenerate.
    
    git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@219344 138bc75d-0d04-0410-961f-82ee72b054a4
---
 ChangeLog    | 6 ++++++
 Makefile.def | 2 +-
 Makefile.in  | 2 +-
 3 files changed, 8 insertions(+), 2 deletions(-)

diff --git ChangeLog ChangeLog
index 325f4cc..9012087 100644
--- ChangeLog
+++ ChangeLog
@@ -1,3 +1,9 @@
+2015-01-08  Thomas Schwinge  <thomas@codesourcery.com>
+
+	* Makefile.def (dependencies) <all-target-liboffloadmic>: Depend on
+	all-target-libgomp.
+	* Makefile.in: Regenerate.
+
 2015-01-06  Eric Botcazou  <ebotcazou@adacore.com>
 
 	* MAINTAINERS (CPU Port Maintainers): Add myself for Visium port.
diff --git Makefile.def Makefile.def
index ffab409..15ab613 100644
--- Makefile.def
+++ Makefile.def
@@ -551,7 +551,7 @@ dependencies = { module=configure-target-libvtv; on=all-target-libstdc++-v3; };
 // generated by the libgomp configure.  Unfortunately, due to the use of
 //  recursive make, we can't be that specific.
 dependencies = { module=all-target-libstdc++-v3; on=configure-target-libgomp; };
-dependencies = { module=all-target-liboffloadmic; on=configure-target-libgomp; };
+dependencies = { module=all-target-liboffloadmic; on=all-target-libgomp; };
 
 dependencies = { module=install-target-libgo; on=install-target-libatomic; };
 dependencies = { module=install-target-libgfortran; on=install-target-libquadmath; };
diff --git Makefile.in Makefile.in
index 7355bf1..428898a 100644
--- Makefile.in
+++ Makefile.in
@@ -48893,7 +48893,7 @@ all-stage3-target-libstdc++-v3: maybe-configure-stage3-target-libgomp
 all-stage4-target-libstdc++-v3: maybe-configure-stage4-target-libgomp
 all-stageprofile-target-libstdc++-v3: maybe-configure-stageprofile-target-libgomp
 all-stagefeedback-target-libstdc++-v3: maybe-configure-stagefeedback-target-libgomp
-all-target-liboffloadmic: maybe-configure-target-libgomp
+all-target-liboffloadmic: maybe-all-target-libgomp
 install-target-libgo: maybe-install-target-libatomic
 install-target-libgfortran: maybe-install-target-libquadmath
 install-target-libgfortran: maybe-install-target-libgcc


Grüße,
 Thomas

[-- Attachment #2: signature.asc --]
[-- Type: application/pgp-signature, Size: 472 bytes --]

Re: [PATCH 3/4] Add libgomp plugin for Intel MIC

[Thomas Schwinge]

[-- Attachment #1: Type: text/plain, Size: 4036 bytes --]

Hi!

On Tue, 21 Oct 2014 21:24:13 +0400, Ilya Verbin <iverbin@gmail.com> wrote:
> This patch contains a plugin for libgomp and appropriate changes for makefiles.
> 
> The plugin uses liboffloadmic_host.so to interact with the device (or with an
> emulator).  Also the patch contains offload_target_main executable, which is the
> corresponding target side part of a libgomp plugin, and it uses
> liboffloadmic_target.so.
> 
> The plugin builds automatically with liboffloadmic.

With recent GCC trunk sources, builds of the Intel MIC Offload Plugin
fail as follows:

    libtool: compile:  [...]/build-gcc/./gcc/xg++ [...] -I[...]/install/offload-x86_64-intelmicemul-linux-gnu/x86_64-intelmicemul-linux-gnu/liboffloadmic/plugin -I[...]/install/offload-x86_64-intelmicemul-linux-gnu/lib/gcc/x86_64-intelmicemul-linux-gnu/6.0.0/include -c [...]/source-gcc/liboffloadmic/plugin/libgomp-plugin-intelmic.cpp  -fPIC -DPIC -o .libs/libgomp_plugin_intelmic_la-libgomp-plugin-intelmic.o
    In file included from [...]/source-gcc/liboffloadmic/plugin/libgomp-plugin-intelmic.cpp:40:0:
    [...]/install/offload-x86_64-intelmicemul-linux-gnu/lib/gcc/x86_64-intelmicemul-linux-gnu/6.0.0/include/main_target_image.h:8628:1: error: narrowing conversion of '192' from 'int' to 'char' inside { } [-Wnarrowing]
     };
     ^
    [...]/install/offload-x86_64-intelmicemul-linux-gnu/lib/gcc/x86_64-intelmicemul-linux-gnu/6.0.0/include/main_target_image.h:8628:1: error: narrowing conversion of '192' from 'int' to 'char' inside { } [-Wnarrowing]
    [...]/install/offload-x86_64-intelmicemul-linux-gnu/lib/gcc/x86_64-intelmicemul-linux-gnu/6.0.0/include/main_target_image.h:8628:1: error: narrowing conversion of '164' from 'int' to 'char' inside { } [-Wnarrowing]
    [many more]

Apart from the actual compilation error, it is surprising for me to see
the GCC build reference/depend on the Intel MIC offloading compiler's
installation directory (which I built and installed earlier),
[...]/install/offload-x86_64-intelmicemul-linux-gnu/ -- is that the
correct thing to do?  Shouldn't the GCC build be self-contained?  (I have
not yet made an attempt to understand how the target and device
liboffloadmic builds work together.)

This main_target_image.h file is coming from here:

> --- /dev/null
> +++ b/liboffloadmic/plugin/Makefile.am
> @@ -0,0 +1,123 @@
> +# Plugin for offload execution on Intel MIC devices.

> +main_target_image.h: offload_target_main
> +	@echo -n "const int image_size = " > $@
> +	@stat -c '%s' $< >> $@
> +	@echo ";" >> $@
> +	@echo "struct MainTargetImage {" >> $@
> +	@echo "  int64_t size;" >> $@
> +	@echo "  char name[sizeof \"offload_target_main\"];" >> $@
> +	@echo "  char data[image_size];" >> $@
> +	@echo "};" >> $@
> +	@echo "extern \"C\" const MainTargetImage main_target_image = {" >> $@
> +	@echo "  image_size, \"offload_target_main\"," >> $@
> +	@cat $< | xxd -include >> $@
> +	@echo "};" >> $@
> +
> +offload_target_main: $(liboffload_dir)/ofldbegin.o offload_target_main.o $(liboffload_dir)/ofldend.o
> +	$(CXX) $(AM_LDFLAGS) $^ -o $@
> +
> +offload_target_main.o: offload_target_main.cpp
> +	$(CXX) $(AM_CXXFLAGS) $(AM_CPPFLAGS) -c $< -o $@

Here, I note that the xxd tool is being used, which in my distribution is
part of the Vim editor's package, which -- as far as I know -- is not
currently declared as a build dependency of GCC?

Anyway, all that aside for the moment -- OK to commit the following?

--- liboffloadmic/plugin/Makefile.am
+++ liboffloadmic/plugin/Makefile.am
@@ -69,7 +69,7 @@ main_target_image.h: offload_target_main
 	@echo "struct MainTargetImage {" >> $@
 	@echo "  int64_t size;" >> $@
 	@echo "  char name[sizeof \"offload_target_main\"];" >> $@
-	@echo "  char data[image_size];" >> $@
+	@echo "  uint8_t data[image_size];" >> $@
 	@echo "};" >> $@
 	@echo "extern \"C\" const MainTargetImage main_target_image = {" >> $@
 	@echo "  image_size, \"offload_target_main\"," >> $@


Grüße,
 Thomas

[-- Attachment #2: signature.asc --]
[-- Type: application/pgp-signature, Size: 472 bytes --]

Re: [PATCH 3/4] Add libgomp plugin for Intel MIC

[Ilya Verbin]

> On 8 июля 2015 г., at 17:16, Thomas Schwinge <thomas@codesourcery.com> wrote:
> 
> Hi!
> 
>> On Tue, 21 Oct 2014 21:24:13 +0400, Ilya Verbin <iverbin@gmail.com> wrote:
>> This patch contains a plugin for libgomp and appropriate changes for makefiles.
>> 
>> The plugin uses liboffloadmic_host.so to interact with the device (or with an
>> emulator).  Also the patch contains offload_target_main executable, which is the
>> corresponding target side part of a libgomp plugin, and it uses
>> liboffloadmic_target.so.
>> 
>> The plugin builds automatically with liboffloadmic.
> 
> With recent GCC trunk sources, builds of the Intel MIC Offload Plugin
> fail as follows:
> 
>    libtool: compile:  [...]/build-gcc/./gcc/xg++ [...] -I[...]/install/offload-x86_64-intelmicemul-linux-gnu/x86_64-intelmicemul-linux-gnu/liboffloadmic/plugin -I[...]/install/offload-x86_64-intelmicemul-linux-gnu/lib/gcc/x86_64-intelmicemul-linux-gnu/6.0.0/include -c [...]/source-gcc/liboffloadmic/plugin/libgomp-plugin-intelmic.cpp  -fPIC -DPIC -o .libs/libgomp_plugin_intelmic_la-libgomp-plugin-intelmic.o
>    In file included from [...]/source-gcc/liboffloadmic/plugin/libgomp-plugin-intelmic.cpp:40:0:
>    [...]/install/offload-x86_64-intelmicemul-linux-gnu/lib/gcc/x86_64-intelmicemul-linux-gnu/6.0.0/include/main_target_image.h:8628:1: error: narrowing conversion of '192' from 'int' to 'char' inside { } [-Wnarrowing]
>     };
>     ^
>    [...]/install/offload-x86_64-intelmicemul-linux-gnu/lib/gcc/x86_64-intelmicemul-linux-gnu/6.0.0/include/main_target_image.h:8628:1: error: narrowing conversion of '192' from 'int' to 'char' inside { } [-Wnarrowing]
>    [...]/install/offload-x86_64-intelmicemul-linux-gnu/lib/gcc/x86_64-intelmicemul-linux-gnu/6.0.0/include/main_target_image.h:8628:1: error: narrowing conversion of '164' from 'int' to 'char' inside { } [-Wnarrowing]
>    [many more]
> 
> Apart from the actual compilation error, it is surprising for me to see
> the GCC build reference/depend on the Intel MIC offloading compiler's
> installation directory (which I built and installed earlier),
> [...]/install/offload-x86_64-intelmicemul-linux-gnu/ -- is that the
> correct thing to do?  Shouldn't the GCC build be self-contained?  (I have
> not yet made an attempt to understand how the target and device
> liboffloadmic builds work together.)
> 
> This main_target_image.h file is coming from here:
> 
>> --- /dev/null
>> +++ b/liboffloadmic/plugin/Makefile.am
>> @@ -0,0 +1,123 @@
>> +# Plugin for offload execution on Intel MIC devices.
> 
>> +main_target_image.h: offload_target_main
>> +    @echo -n "const int image_size = " > $@
>> +    @stat -c '%s' $< >> $@
>> +    @echo ";" >> $@
>> +    @echo "struct MainTargetImage {" >> $@
>> +    @echo "  int64_t size;" >> $@
>> +    @echo "  char name[sizeof \"offload_target_main\"];" >> $@
>> +    @echo "  char data[image_size];" >> $@
>> +    @echo "};" >> $@
>> +    @echo "extern \"C\" const MainTargetImage main_target_image = {" >> $@
>> +    @echo "  image_size, \"offload_target_main\"," >> $@
>> +    @cat $< | xxd -include >> $@
>> +    @echo "};" >> $@
>> +
>> +offload_target_main: $(liboffload_dir)/ofldbegin.o offload_target_main.o $(liboffload_dir)/ofldend.o
>> +    $(CXX) $(AM_LDFLAGS) $^ -o $@
>> +
>> +offload_target_main.o: offload_target_main.cpp
>> +    $(CXX) $(AM_CXXFLAGS) $(AM_CPPFLAGS) -c $< -o $@
> 
> Here, I note that the xxd tool is being used, which in my distribution is
> part of the Vim editor's package, which -- as far as I know -- is not
> currently declared as a build dependency of GCC?
> 
> Anyway, all that aside for the moment -- OK to commit the following?
> 
> --- liboffloadmic/plugin/Makefile.am
> +++ liboffloadmic/plugin/Makefile.am
> @@ -69,7 +69,7 @@ main_target_image.h: offload_target_main
>    @echo "struct MainTargetImage {" >> $@
>    @echo "  int64_t size;" >> $@
>    @echo "  char name[sizeof \"offload_target_main\"];" >> $@
> -    @echo "  char data[image_size];" >> $@
> +    @echo "  uint8_t data[image_size];" >> $@
>    @echo "};" >> $@
>    @echo "extern \"C\" const MainTargetImage main_target_image = {" >> $@
>    @echo "  image_size, \"offload_target_main\"," >> $@
> 
> 
> Grüße,
> Thomas

Ok to me, thanks.

The plugin consists of 2 parts: offload_target_main is a target part, which is embedded into the host part (libgomp plugin itself). Target part is linked with liboffloadmic_target.so and host part is linked with liboffloadmic_host.so. Both offload_target_main and liboffloadmic_target.so are compiled by the target compiler during its build.

As for xxd, I've found its usage in some Makefile inside gcc tree, so I thought it's ok to use it.

  -- Ilya

Re: [PATCH 3/4] Add libgomp plugin for Intel MIC

[Thomas Schwinge]

[-- Attachment #1: Type: text/plain, Size: 4302 bytes --]

Hi!

On Wed, 8 Jul 2015 18:13:56 +0300, Ilya Verbin <iverbin@gmail.com> wrote:
> 
> > On 8 июля 2015 г., at 17:16, Thomas Schwinge <thomas@codesourcery.com> wrote:
> > With recent GCC trunk sources, builds of the Intel MIC Offload Plugin
> > fail as follows: [...]

> > [...] -- OK to commit the following?

> Ok to me, thanks.

Committed in r225562:

commit cacef506e4205bac13a0dd1de238d1a8cc78af28
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date:   Wed Jul 8 15:47:59 2015 +0000

    liboffloadmic plugin: Address -Wnarrowing diagnostics
    
        libtool: compile:  [...]/build-gcc/./gcc/xg++ [...] -c [...]/source-gcc/liboffloadmic/plugin/libgomp-plugin-intelmic.cpp [...]
        In file included from [...]/source-gcc/liboffloadmic/plugin/libgomp-plugin-intelmic.cpp:40:0:
        [...]/install/offload-x86_64-intelmicemul-linux-gnu/lib/gcc/x86_64-intelmicemul-linux-gnu/6.0.0/include/main_target_image.h:8628:1: error: narrowing conversion of '192' from 'int' to 'char' inside { } [-Wnarrowing]
         };
         ^
        [...]/install/offload-x86_64-intelmicemul-linux-gnu/lib/gcc/x86_64-intelmicemul-linux-gnu/6.0.0/include/main_target_image.h:8628:1: error: narrowing conversion of '192' from 'int' to 'char' inside { } [-Wnarrowing]
        [...]/install/offload-x86_64-intelmicemul-linux-gnu/lib/gcc/x86_64-intelmicemul-linux-gnu/6.0.0/include/main_target_image.h:8628:1: error: narrowing conversion of '164' from 'int' to 'char' inside { } [-Wnarrowing]
        [many more]
    
    	liboffloadmic/
    	* plugin/Makefile.am (main_target_image.h): Change type of data
    	member in struct MainTargetImage to uint8_t.
    	* plugin/Makefile.in: Regenerate.
    
    git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@225562 138bc75d-0d04-0410-961f-82ee72b054a4
---
 liboffloadmic/ChangeLog          |    6 ++++++
 liboffloadmic/plugin/Makefile.am |    2 +-
 liboffloadmic/plugin/Makefile.in |    2 +-
 3 files changed, 8 insertions(+), 2 deletions(-)

diff --git liboffloadmic/ChangeLog liboffloadmic/ChangeLog
index 01fb9f4..b0f9e90 100644
--- liboffloadmic/ChangeLog
+++ liboffloadmic/ChangeLog
@@ -1,3 +1,9 @@
+2015-07-08  Thomas Schwinge  <thomas@codesourcery.com>
+
+	* plugin/Makefile.am (main_target_image.h): Change type of data
+	member of struct MainTargetImage to uint8_t.
+	* plugin/Makefile.in: Regenerate.
+
 2015-05-13  Michael Haubenwallner  <michael.haubenwallner@ssi-schaefer.com>
 
 	* Makefile.in: Regenerated with automake-1.11.6.
diff --git liboffloadmic/plugin/Makefile.am liboffloadmic/plugin/Makefile.am
index a814f0c..19d69ab 100644
--- liboffloadmic/plugin/Makefile.am
+++ liboffloadmic/plugin/Makefile.am
@@ -69,7 +69,7 @@ main_target_image.h: offload_target_main
 	@echo "struct MainTargetImage {" >> $@
 	@echo "  int64_t size;" >> $@
 	@echo "  char name[sizeof \"offload_target_main\"];" >> $@
-	@echo "  char data[image_size];" >> $@
+	@echo "  uint8_t data[image_size];" >> $@
 	@echo "};" >> $@
 	@echo "extern \"C\" const MainTargetImage main_target_image = {" >> $@
 	@echo "  image_size, \"offload_target_main\"," >> $@
diff --git liboffloadmic/plugin/Makefile.in liboffloadmic/plugin/Makefile.in
index 6f7eec9..19a1a96 100644
--- liboffloadmic/plugin/Makefile.in
+++ liboffloadmic/plugin/Makefile.in
@@ -715,7 +715,7 @@ main_target_image.h: offload_target_main
 	@echo "struct MainTargetImage {" >> $@
 	@echo "  int64_t size;" >> $@
 	@echo "  char name[sizeof \"offload_target_main\"];" >> $@
-	@echo "  char data[image_size];" >> $@
+	@echo "  uint8_t data[image_size];" >> $@
 	@echo "};" >> $@
 	@echo "extern \"C\" const MainTargetImage main_target_image = {" >> $@
 	@echo "  image_size, \"offload_target_main\"," >> $@


Thanks for the explanation:

> The plugin consists of 2 parts: offload_target_main is a target part, which is embedded into the host part (libgomp plugin itself). Target part is linked with liboffloadmic_target.so and host part is linked with liboffloadmic_host.so. Both offload_target_main and liboffloadmic_target.so are compiled by the target compiler during its build.
> 
> As for xxd, I've found its usage in some Makefile inside gcc tree, so I thought it's ok to use it.


Grüße,
 Thomas

[-- Attachment #2: Type: application/pgp-signature, Size: 472 bytes --]

Re: [PATCH 3/4] Add libgomp plugin for Intel MIC

[Ilya Verbin]

On Wed, Jul 08, 2015 at 16:16:44 +0200, Thomas Schwinge wrote:
> > --- /dev/null
> > +++ b/liboffloadmic/plugin/Makefile.am
> > @@ -0,0 +1,123 @@
> > +# Plugin for offload execution on Intel MIC devices.
> 
> > +main_target_image.h: offload_target_main
> > +	@echo -n "const int image_size = " > $@
> > +	@stat -c '%s' $< >> $@
> > +	@echo ";" >> $@
> > +	@echo "struct MainTargetImage {" >> $@
> > +	@echo "  int64_t size;" >> $@
> > +	@echo "  char name[sizeof \"offload_target_main\"];" >> $@
> > +	@echo "  char data[image_size];" >> $@
> > +	@echo "};" >> $@
> > +	@echo "extern \"C\" const MainTargetImage main_target_image = {" >> $@
> > +	@echo "  image_size, \"offload_target_main\"," >> $@
> > +	@cat $< | xxd -include >> $@
> > +	@echo "};" >> $@
> > +
> > +offload_target_main: $(liboffload_dir)/ofldbegin.o offload_target_main.o $(liboffload_dir)/ofldend.o
> > +	$(CXX) $(AM_LDFLAGS) $^ -o $@
> > +
> > +offload_target_main.o: offload_target_main.cpp
> > +	$(CXX) $(AM_CXXFLAGS) $(AM_CPPFLAGS) -c $< -o $@
> 
> Here, I note that the xxd tool is being used, which in my distribution is
> part of the Vim editor's package, which -- as far as I know -- is not
> currently declared as a build dependency of GCC?

We have a patch, which checks for xxd availability, is it ok for trunk?


2015-07-23  Maxim Blumenthal  <maxim.blumenthal@intel.com>

	* configure.ac: Add a check for xxd presence when the target is
	intelmic or intelmicemul.
	* configure: Regenerate.


diff --git a/configure b/configure
index 5ba9489..bd8fed8 100755
--- a/configure
+++ b/configure
[regenerate]

diff --git a/configure.ac b/configure.ac
index 2ff9be0..63eebfc 100644
--- a/configure.ac
+++ b/configure.ac
@@ -494,6 +494,17 @@ else
 fi])
 AC_SUBST(extra_liboffloadmic_configure_flags)
 
+# Check if xxd is present in the system
+# when the target is intelmic or intelmicemul.
+case "${target}" in
+  *-intelmic-* | *-intelmicemul-*)
+    AC_CHECK_PROG(xxd_present, xxd, "yes", "no")
+    if test "$xxd_present" = "no"; then
+      AC_MSG_ERROR([cannot find xxd])
+    fi
+    ;;
+esac
+
 # Save it here so that, even in case of --enable-libgcj, if the Java
 # front-end isn't enabled, we still get libgcj disabled.
 libgcj_saved=$libgcj


  -- Ilya

Re: [PATCH 3/4] Add libgomp plugin for Intel MIC

[Jakub Jelinek]

On Thu, Jul 23, 2015 at 09:50:55PM +0300, Ilya Verbin wrote:
> > Here, I note that the xxd tool is being used, which in my distribution is
> > part of the Vim editor's package, which -- as far as I know -- is not
> > currently declared as a build dependency of GCC?
> 
> We have a patch, which checks for xxd availability, is it ok for trunk?

I'd prefer at least some alternatives.  E.g. the following xxd.py

#!/usr/bin/python
import sys
with open(sys.argv[1],"rb") as f:
    nextblock = f.read(12)
    while 1:
	block = nextblock
	nextblock = f.read(12)
	if block == "":
	    break
	str = ""
	for ch in block:
	    if str == "":
		str = "  "
	    else:
		str += ", "
	    if ord(ch) < 10:
		str += "0x0" + chr(ord('0')+ord(ch))
	    elif ord(ch) < 16:
		str += "0x0" + chr(ord('a')+ord(ch)-10)
	    else:
		str += hex(ord(ch))
	if nextblock != "":
	    str += ","
	print str

	python ./xxd.py $< >> $@
does the same thing as
	cat $< | xxd -include >> $@
(CCing David as python expert, my python knowledge is limited and
15 years old, not sure how portable this is (python 2 vs. python 3, and
even python 2 minimal versions)).

Thus, perhaps configure could check for python that can handle this,
or xxd, and substitute the right command into the makefile and
only bail out if neither is found?

	Jakub

Re: [PATCH 3/4] Add libgomp plugin for Intel MIC

[David Malcolm]

[-- Attachment #1: Type: text/plain, Size: 1426 bytes --]

On Fri, 2015-07-24 at 10:01 +0200, Jakub Jelinek wrote:
> #!/usr/bin/python
> import sys
> with open(sys.argv[1],"rb") as f:
>     nextblock = f.read(12)
>     while 1:
>         block = nextblock
>         nextblock = f.read(12)
>         if block == "":
>             break
>         str = ""
>         for ch in block:
>             if str == "":
>                 str = "  "
>             else:
>                 str += ", "
>             if ord(ch) < 10:
>                 str += "0x0" + chr(ord('0')+ord(ch))
>             elif ord(ch) < 16:
>                 str += "0x0" + chr(ord('a')+ord(ch)-10)
>             else:
>                 str += hex(ord(ch))
>         if nextblock != "":
>             str += ","
>         print str
> 
>         python ./xxd.py $< >> $@
> does the same thing as
>         cat $< | xxd -include >> $@
> (CCing David as python expert, my python knowledge is limited and
> 15 years old, not sure how portable this is (python 2 vs. python 3,
> and
> even python 2 minimal versions)).

It doesn't work with Python 3 for various reasons ("print" syntax, and
str vs bytes issues).

I'm attaching a version which works with both Python 2 and Python 3
(2.7.5 and 3.3.2 were the versions I tried).

It ought to work with much older python 2 versions (as your script
appears to), but I don't have them handy.

Presumably it would need a license header and some descriptive comments.

(snip)

Dave

[-- Attachment #2: xxd.py --]
[-- Type: text/x-python, Size: 1019 bytes --]

#!/usr/bin/python
import sys

if sys.version_info[0] == 2:
    # Python 2:
    # 'block' below is an instance of str; iterating over it gives us
    # str instances of len 1.
    def get_byte(ch):
        return ord(ch)
else:
    # Python 3:
    # 'block' below is an instance of bytes; iterating over it gives us
    # instances of int, in the range 0-255.
    def get_byte(ch):
        return ch

with open(sys.argv[1],"rb") as f:
    nextblock = f.read(12)
    while 1:
        block = nextblock
        nextblock = f.read(12)
        if not block:
            break
        str = ""
        for item in block:
            byte = get_byte(item)
            if str == "":
                str = "  "
            else:
                str += ", "
            if byte < 10:
                str += "0x0" + chr(ord('0')+byte)
            elif byte < 16:
                str += "0x0" + chr(ord('a')+byte-10)
            else:
                str += hex(byte)
        if nextblock:
            str += ","
        print(str)

Re: [PATCH 3/4] Add libgomp plugin for Intel MIC

[Maxim Blumental]

[-- Attachment #1: Type: text/plain, Size: 1678 bytes --]

 Applied the idea with python script alternative. Review, please.

2015-07-24 17:18 GMT+03:00 David Malcolm <dmalcolm@redhat.com>:
> On Fri, 2015-07-24 at 10:01 +0200, Jakub Jelinek wrote:
>> #!/usr/bin/python
>> import sys
>> with open(sys.argv[1],"rb") as f:
>>     nextblock = f.read(12)
>>     while 1:
>>         block = nextblock
>>         nextblock = f.read(12)
>>         if block == "":
>>             break
>>         str = ""
>>         for ch in block:
>>             if str == "":
>>                 str = "  "
>>             else:
>>                 str += ", "
>>             if ord(ch) < 10:
>>                 str += "0x0" + chr(ord('0')+ord(ch))
>>             elif ord(ch) < 16:
>>                 str += "0x0" + chr(ord('a')+ord(ch)-10)
>>             else:
>>                 str += hex(ord(ch))
>>         if nextblock != "":
>>             str += ","
>>         print str
>>
>>         python ./xxd.py $< >> $@
>> does the same thing as
>>         cat $< | xxd -include >> $@
>> (CCing David as python expert, my python knowledge is limited and
>> 15 years old, not sure how portable this is (python 2 vs. python 3,
>> and
>> even python 2 minimal versions)).
>
> It doesn't work with Python 3 for various reasons ("print" syntax, and
> str vs bytes issues).
>
> I'm attaching a version which works with both Python 2 and Python 3
> (2.7.5 and 3.3.2 were the versions I tried).
>
> It ought to work with much older python 2 versions (as your script
> appears to), but I don't have them handy.
>
> Presumably it would need a license header and some descriptive comments.
>
> (snip)
>
> Dave



-- 


---------------------
Sincerely yours,
Maxim Blumental

[-- Attachment #2: ChangeLog.txt --]
[-- Type: text/plain, Size: 489 bytes --]

2015-07-28  Maxim Blumenthal  <maxim.blumenthal@intel.com>

	* configure.ac: Add a check for xxd or python presence when the target
	is intelmic or intelmicemul.
	* configure: Regenerate.
	* liboffloadmic/plugin/Makefile.am: Add a condition into
	make_target_image.h generating code.  This condition performs an
	action with either xxd or a special python script during the
	generating.
	* liboffloadmic/plugin/xxd.py: New file.
	* liboffloadmic/plugin/Makefile.in: Regenerate.

[-- Attachment #3: xxd_check.patch --]
[-- Type: application/octet-stream, Size: 8208 bytes --]

commit ab7826ee4ee5f2920b015d01aa538c3ebf50dccd
Author: Maxim Blumenthal <maxim.blumenthal@intel.com>
Date:   Mon Jul 27 20:15:01 2015 +0300

    Check xxd

diff --git a/configure b/configure
index 6d7152e..faad566 100755
--- a/configure
+++ b/configure
@@ -674,6 +674,8 @@ LDFLAGS
 CFLAGS
 CC
 EXTRA_CONFIGARGS_LIBJAVA
+python_present
+xxd_present
 extra_liboffloadmic_configure_flags
 target_subdir
 host_subdir
@@ -3131,6 +3133,92 @@ fi
 
 
 
+# Check if xxd is present in the system
+# when the target is intelmic or intelmicemul.
+case "${target}" in
+  *-intelmic-* | *-intelmicemul-*)
+    # Extract the first word of "xxd", so it can be a program name with args.
+set dummy xxd; ac_word=$2
+{ $as_echo "$as_me:${as_lineno-$LINENO}: checking for $ac_word" >&5
+$as_echo_n "checking for $ac_word... " >&6; }
+if test "${ac_cv_prog_xxd_present+set}" = set; then :
+  $as_echo_n "(cached) " >&6
+else
+  if test -n "$xxd_present"; then
+  ac_cv_prog_xxd_present="$xxd_present" # Let the user override the test.
+else
+as_save_IFS=$IFS; IFS=$PATH_SEPARATOR
+for as_dir in $PATH
+do
+  IFS=$as_save_IFS
+  test -z "$as_dir" && as_dir=.
+    for ac_exec_ext in '' $ac_executable_extensions; do
+  if { test -f "$as_dir/$ac_word$ac_exec_ext" && $as_test_x "$as_dir/$ac_word$ac_exec_ext"; }; then
+    ac_cv_prog_xxd_present=""yes""
+    $as_echo "$as_me:${as_lineno-$LINENO}: found $as_dir/$ac_word$ac_exec_ext" >&5
+    break 2
+  fi
+done
+  done
+IFS=$as_save_IFS
+
+  test -z "$ac_cv_prog_xxd_present" && ac_cv_prog_xxd_present=""no""
+fi
+fi
+xxd_present=$ac_cv_prog_xxd_present
+if test -n "$xxd_present"; then
+  { $as_echo "$as_me:${as_lineno-$LINENO}: result: $xxd_present" >&5
+$as_echo "$xxd_present" >&6; }
+else
+  { $as_echo "$as_me:${as_lineno-$LINENO}: result: no" >&5
+$as_echo "no" >&6; }
+fi
+
+
+    # Extract the first word of "python", so it can be a program name with args.
+set dummy python; ac_word=$2
+{ $as_echo "$as_me:${as_lineno-$LINENO}: checking for $ac_word" >&5
+$as_echo_n "checking for $ac_word... " >&6; }
+if test "${ac_cv_prog_python_present+set}" = set; then :
+  $as_echo_n "(cached) " >&6
+else
+  if test -n "$python_present"; then
+  ac_cv_prog_python_present="$python_present" # Let the user override the test.
+else
+as_save_IFS=$IFS; IFS=$PATH_SEPARATOR
+for as_dir in $PATH
+do
+  IFS=$as_save_IFS
+  test -z "$as_dir" && as_dir=.
+    for ac_exec_ext in '' $ac_executable_extensions; do
+  if { test -f "$as_dir/$ac_word$ac_exec_ext" && $as_test_x "$as_dir/$ac_word$ac_exec_ext"; }; then
+    ac_cv_prog_python_present=""yes""
+    $as_echo "$as_me:${as_lineno-$LINENO}: found $as_dir/$ac_word$ac_exec_ext" >&5
+    break 2
+  fi
+done
+  done
+IFS=$as_save_IFS
+
+  test -z "$ac_cv_prog_python_present" && ac_cv_prog_python_present=""no""
+fi
+fi
+python_present=$ac_cv_prog_python_present
+if test -n "$python_present"; then
+  { $as_echo "$as_me:${as_lineno-$LINENO}: result: $python_present" >&5
+$as_echo "$python_present" >&6; }
+else
+  { $as_echo "$as_me:${as_lineno-$LINENO}: result: no" >&5
+$as_echo "no" >&6; }
+fi
+
+
+    if test "$xxd_present$python_present" = "nono"; then
+      as_fn_error "cannot find neither xxd nor python" "$LINENO" 5
+    fi
+    ;;
+esac
+
 # Save it here so that, even in case of --enable-libgcj, if the Java
 # front-end isn't enabled, we still get libgcj disabled.
 libgcj_saved=$libgcj
diff --git a/configure.ac b/configure.ac
index fbc49ce..f429a1f 100644
--- a/configure.ac
+++ b/configure.ac
@@ -494,6 +494,18 @@ else
 fi])
 AC_SUBST(extra_liboffloadmic_configure_flags)
 
+# Check if xxd is present in the system
+# when the target is intelmic or intelmicemul.
+case "${target}" in
+  *-intelmic-* | *-intelmicemul-*)
+    AC_CHECK_PROG(xxd_present, xxd, "yes", "no")
+    AC_CHECK_PROG(python_present, python, "yes", "no")
+    if test "$xxd_present$python_present" = "nono"; then
+      AC_MSG_ERROR([cannot find neither xxd nor python])
+    fi
+    ;;
+esac
+
 # Save it here so that, even in case of --enable-libgcj, if the Java
 # front-end isn't enabled, we still get libgcj disabled.
 libgcj_saved=$libgcj
diff --git a/liboffloadmic/plugin/Makefile.am b/liboffloadmic/plugin/Makefile.am
index 19d69ab..cb93309 100644
--- a/liboffloadmic/plugin/Makefile.am
+++ b/liboffloadmic/plugin/Makefile.am
@@ -49,6 +49,10 @@ target_prefix_dir = $(libdir)/gcc/$(accel_target)/$(gcc_version)$(MULTISUBDIR)
 target_build_dir = $(accel_search_dir)/$(accel_target)$(MULTISUBDIR)/liboffloadmic/plugin
 target_install_dir = $(accel_search_dir)/lib/gcc/$(accel_target)/$(gcc_version)$(MULTISUBDIR)
 
+XXD_PY = $(source_dir)/xxd.py
+xxd_path=`which xxd`
+python_path=`which python`
+
 if PLUGIN_HOST
   toolexeclib_LTLIBRARIES = libgomp-plugin-intelmic.la
   libgomp_plugin_intelmic_la_SOURCES = libgomp-plugin-intelmic.cpp
@@ -73,7 +77,7 @@ main_target_image.h: offload_target_main
 	@echo "};" >> $@
 	@echo "extern \"C\" const MainTargetImage main_target_image = {" >> $@
 	@echo "  image_size, \"offload_target_main\"," >> $@
-	@cat $< | xxd -include >> $@
+	@if test "x$(xxd_path)" != "x"; then cat $< | $(xxd_path) -include >> $@; else $(python_path) $(XXD_PY) $< >> $@; fi;
 	@echo "};" >> $@
 
 offload_target_main: $(liboffload_dir)/ofldbegin.o offload_target_main.o $(liboffload_dir)/ofldend.o
diff --git a/liboffloadmic/plugin/Makefile.in b/liboffloadmic/plugin/Makefile.in
index 19a1a96..12fe5ca 100644
--- a/liboffloadmic/plugin/Makefile.in
+++ b/liboffloadmic/plugin/Makefile.in
@@ -318,6 +318,9 @@ libsubincludedir = $(libdir)/gcc/$(target_alias)/$(gcc_version)$(MULTISUBDIR)/in
 target_prefix_dir = $(libdir)/gcc/$(accel_target)/$(gcc_version)$(MULTISUBDIR)
 target_build_dir = $(accel_search_dir)/$(accel_target)$(MULTISUBDIR)/liboffloadmic/plugin
 target_install_dir = $(accel_search_dir)/lib/gcc/$(accel_target)/$(gcc_version)$(MULTISUBDIR)
+XXD_PY = $(source_dir)/xxd.py
+xxd_path = `which xxd`
+python_path = `which python`
 @PLUGIN_HOST_TRUE@toolexeclib_LTLIBRARIES = libgomp-plugin-intelmic.la
 @PLUGIN_HOST_TRUE@libgomp_plugin_intelmic_la_SOURCES = libgomp-plugin-intelmic.cpp
 @PLUGIN_HOST_TRUE@libgomp_plugin_intelmic_la_CPPFLAGS = $(CPPFLAGS) -DLINUX -DCOI_LIBRARY_VERSION=2 -DMYO_SUPPORT -DOFFLOAD_DEBUG=1 -DSEP_SUPPORT -DTIMING_SUPPORT -DHOST_LIBRARY=1 -I$(coi_inc_dir) -I$(myo_inc_dir) -I$(liboffload_src_dir) -I$(libgomp_src_dir) -I$(libgomp_dir) -I$(target_prefix_dir)/include -I$(target_build_dir) -I$(target_install_dir)/include
@@ -719,7 +722,7 @@ main_target_image.h: offload_target_main
 	@echo "};" >> $@
 	@echo "extern \"C\" const MainTargetImage main_target_image = {" >> $@
 	@echo "  image_size, \"offload_target_main\"," >> $@
-	@cat $< | xxd -include >> $@
+	@if test "x$(xxd_path)" != "x"; then cat $< | $(xxd_path) -include >> $@; else $(python_path) $(XXD_PY) $< >> $@; fi;
 	@echo "};" >> $@
 
 offload_target_main: $(liboffload_dir)/ofldbegin.o offload_target_main.o $(liboffload_dir)/ofldend.o
diff --git a/liboffloadmic/plugin/xxd.py b/liboffloadmic/plugin/xxd.py
new file mode 100644
index 0000000..fd9d7f0
--- /dev/null
+++ b/liboffloadmic/plugin/xxd.py
@@ -0,0 +1,39 @@
+#!/usr/bin/python
+import sys
+
+if sys.version_info[0] == 2:
+    # Python 2:
+    # 'block' below is an instance of str; iterating over it gives us
+    # str instances of len 1.
+    def get_byte(ch):
+        return ord(ch)
+else:
+    # Python 3:
+    # 'block' below is an instance of bytes; iterating over it gives us
+    # instances of int, in the range 0-255.
+    def get_byte(ch):
+        return ch
+
+with open(sys.argv[1],"rb") as f:
+    nextblock = f.read(12)
+    while 1:
+        block = nextblock
+        nextblock = f.read(12)
+        if not block:
+            break
+        str = ""
+        for item in block:
+            byte = get_byte(item)
+            if str == "":
+                str = "  "
+            else:
+                str += ", "
+            if byte < 10:
+                str += "0x0" + chr(ord('0')+byte)
+            elif byte < 16:
+                str += "0x0" + chr(ord('a')+byte-10)
+            else:
+                str += hex(byte)
+        if nextblock:
+            str += ","
+        print(str)

Re: [PATCH 3/4] Add libgomp plugin for Intel MIC

[Maxim Blumental]

Could you probably review the patch, please?

2015-07-28 18:42 GMT+03:00 Maxim Blumental <bvmaks@gmail.com>:
>  Applied the idea with python script alternative. Review, please.
>
> 2015-07-24 17:18 GMT+03:00 David Malcolm <dmalcolm@redhat.com>:
>> On Fri, 2015-07-24 at 10:01 +0200, Jakub Jelinek wrote:
>>> #!/usr/bin/python
>>> import sys
>>> with open(sys.argv[1],"rb") as f:
>>>     nextblock = f.read(12)
>>>     while 1:
>>>         block = nextblock
>>>         nextblock = f.read(12)
>>>         if block == "":
>>>             break
>>>         str = ""
>>>         for ch in block:
>>>             if str == "":
>>>                 str = "  "
>>>             else:
>>>                 str += ", "
>>>             if ord(ch) < 10:
>>>                 str += "0x0" + chr(ord('0')+ord(ch))
>>>             elif ord(ch) < 16:
>>>                 str += "0x0" + chr(ord('a')+ord(ch)-10)
>>>             else:
>>>                 str += hex(ord(ch))
>>>         if nextblock != "":
>>>             str += ","
>>>         print str
>>>
>>>         python ./xxd.py $< >> $@
>>> does the same thing as
>>>         cat $< | xxd -include >> $@
>>> (CCing David as python expert, my python knowledge is limited and
>>> 15 years old, not sure how portable this is (python 2 vs. python 3,
>>> and
>>> even python 2 minimal versions)).
>>
>> It doesn't work with Python 3 for various reasons ("print" syntax, and
>> str vs bytes issues).
>>
>> I'm attaching a version which works with both Python 2 and Python 3
>> (2.7.5 and 3.3.2 were the versions I tried).
>>
>> It ought to work with much older python 2 versions (as your script
>> appears to), but I don't have them handy.
>>
>> Presumably it would need a license header and some descriptive comments.
>>
>> (snip)
>>
>> Dave
>
>
>
> --
>
>
> ---------------------
> Sincerely yours,
> Maxim Blumental



-- 


---------------------
Sincerely yours,
Maxim Blumental

Re: [PATCH 3/4] Add libgomp plugin for Intel MIC

[David Malcolm]

On Mon, 2015-08-03 at 13:23 +0300, Maxim Blumental wrote:
> Could you probably review the patch, please?

Sorry, I'm not the best person to review the patch: Jakub CCed me for my
knowledge of python, so I ported his script to work with both python 2
and 3, and it ought to work with early python 2 versions (or be easily
fixable).

It looks like you're using the resulting python script I wrote.  Other
than that, I don't have reviewer-level expertise in the domains of the
rest of the patch (e.g. Intel MIC, and or the build system).

> 2015-07-28 18:42 GMT+03:00 Maxim Blumental <bvmaks@gmail.com>:
> >  Applied the idea with python script alternative. Review, please.
> >
> > 2015-07-24 17:18 GMT+03:00 David Malcolm <dmalcolm@redhat.com>:
> >> On Fri, 2015-07-24 at 10:01 +0200, Jakub Jelinek wrote:
> >>> #!/usr/bin/python
> >>> import sys
> >>> with open(sys.argv[1],"rb") as f:
> >>>     nextblock = f.read(12)
> >>>     while 1:
> >>>         block = nextblock
> >>>         nextblock = f.read(12)
> >>>         if block == "":
> >>>             break
> >>>         str = ""
> >>>         for ch in block:
> >>>             if str == "":
> >>>                 str = "  "
> >>>             else:
> >>>                 str += ", "
> >>>             if ord(ch) < 10:
> >>>                 str += "0x0" + chr(ord('0')+ord(ch))
> >>>             elif ord(ch) < 16:
> >>>                 str += "0x0" + chr(ord('a')+ord(ch)-10)
> >>>             else:
> >>>                 str += hex(ord(ch))
> >>>         if nextblock != "":
> >>>             str += ","
> >>>         print str
> >>>
> >>>         python ./xxd.py $< >> $@
> >>> does the same thing as
> >>>         cat $< | xxd -include >> $@
> >>> (CCing David as python expert, my python knowledge is limited and
> >>> 15 years old, not sure how portable this is (python 2 vs. python 3,
> >>> and
> >>> even python 2 minimal versions)).
> >>
> >> It doesn't work with Python 3 for various reasons ("print" syntax, and
> >> str vs bytes issues).
> >>
> >> I'm attaching a version which works with both Python 2 and Python 3
> >> (2.7.5 and 3.3.2 were the versions I tried).
> >>
> >> It ought to work with much older python 2 versions (as your script
> >> appears to), but I don't have them handy.
> >>
> >> Presumably it would need a license header and some descriptive comments.
> >>
> >> (snip)
> >>
> >> Dave

Fwd: [PATCH 3/4] Add libgomp plugin for Intel MIC

[Maxim Blumental]

[-- Attachment #1: Type: text/plain, Size: 66 bytes --]

 Applied the idea with python script alternative. Review, please.

[-- Attachment #2: ChangeLog.txt --]
[-- Type: text/plain, Size: 489 bytes --]

2015-07-28  Maxim Blumenthal  <maxim.blumenthal@intel.com>

	* configure.ac: Add a check for xxd or python presence when the target
	is intelmic or intelmicemul.
	* configure: Regenerate.
	* liboffloadmic/plugin/Makefile.am: Add a condition into
	make_target_image.h generating code.  This condition performs an
	action with either xxd or a special python script during the
	generating.
	* liboffloadmic/plugin/xxd.py: New file.
	* liboffloadmic/plugin/Makefile.in: Regenerate.

[-- Attachment #3: xxd_check.patch --]
[-- Type: application/octet-stream, Size: 8208 bytes --]

commit ab7826ee4ee5f2920b015d01aa538c3ebf50dccd
Author: Maxim Blumenthal <maxim.blumenthal@intel.com>
Date:   Mon Jul 27 20:15:01 2015 +0300

    Check xxd

diff --git a/configure b/configure
index 6d7152e..faad566 100755
--- a/configure
+++ b/configure
@@ -674,6 +674,8 @@ LDFLAGS
 CFLAGS
 CC
 EXTRA_CONFIGARGS_LIBJAVA
+python_present
+xxd_present
 extra_liboffloadmic_configure_flags
 target_subdir
 host_subdir
@@ -3131,6 +3133,92 @@ fi
 
 
 
+# Check if xxd is present in the system
+# when the target is intelmic or intelmicemul.
+case "${target}" in
+  *-intelmic-* | *-intelmicemul-*)
+    # Extract the first word of "xxd", so it can be a program name with args.
+set dummy xxd; ac_word=$2
+{ $as_echo "$as_me:${as_lineno-$LINENO}: checking for $ac_word" >&5
+$as_echo_n "checking for $ac_word... " >&6; }
+if test "${ac_cv_prog_xxd_present+set}" = set; then :
+  $as_echo_n "(cached) " >&6
+else
+  if test -n "$xxd_present"; then
+  ac_cv_prog_xxd_present="$xxd_present" # Let the user override the test.
+else
+as_save_IFS=$IFS; IFS=$PATH_SEPARATOR
+for as_dir in $PATH
+do
+  IFS=$as_save_IFS
+  test -z "$as_dir" && as_dir=.
+    for ac_exec_ext in '' $ac_executable_extensions; do
+  if { test -f "$as_dir/$ac_word$ac_exec_ext" && $as_test_x "$as_dir/$ac_word$ac_exec_ext"; }; then
+    ac_cv_prog_xxd_present=""yes""
+    $as_echo "$as_me:${as_lineno-$LINENO}: found $as_dir/$ac_word$ac_exec_ext" >&5
+    break 2
+  fi
+done
+  done
+IFS=$as_save_IFS
+
+  test -z "$ac_cv_prog_xxd_present" && ac_cv_prog_xxd_present=""no""
+fi
+fi
+xxd_present=$ac_cv_prog_xxd_present
+if test -n "$xxd_present"; then
+  { $as_echo "$as_me:${as_lineno-$LINENO}: result: $xxd_present" >&5
+$as_echo "$xxd_present" >&6; }
+else
+  { $as_echo "$as_me:${as_lineno-$LINENO}: result: no" >&5
+$as_echo "no" >&6; }
+fi
+
+
+    # Extract the first word of "python", so it can be a program name with args.
+set dummy python; ac_word=$2
+{ $as_echo "$as_me:${as_lineno-$LINENO}: checking for $ac_word" >&5
+$as_echo_n "checking for $ac_word... " >&6; }
+if test "${ac_cv_prog_python_present+set}" = set; then :
+  $as_echo_n "(cached) " >&6
+else
+  if test -n "$python_present"; then
+  ac_cv_prog_python_present="$python_present" # Let the user override the test.
+else
+as_save_IFS=$IFS; IFS=$PATH_SEPARATOR
+for as_dir in $PATH
+do
+  IFS=$as_save_IFS
+  test -z "$as_dir" && as_dir=.
+    for ac_exec_ext in '' $ac_executable_extensions; do
+  if { test -f "$as_dir/$ac_word$ac_exec_ext" && $as_test_x "$as_dir/$ac_word$ac_exec_ext"; }; then
+    ac_cv_prog_python_present=""yes""
+    $as_echo "$as_me:${as_lineno-$LINENO}: found $as_dir/$ac_word$ac_exec_ext" >&5
+    break 2
+  fi
+done
+  done
+IFS=$as_save_IFS
+
+  test -z "$ac_cv_prog_python_present" && ac_cv_prog_python_present=""no""
+fi
+fi
+python_present=$ac_cv_prog_python_present
+if test -n "$python_present"; then
+  { $as_echo "$as_me:${as_lineno-$LINENO}: result: $python_present" >&5
+$as_echo "$python_present" >&6; }
+else
+  { $as_echo "$as_me:${as_lineno-$LINENO}: result: no" >&5
+$as_echo "no" >&6; }
+fi
+
+
+    if test "$xxd_present$python_present" = "nono"; then
+      as_fn_error "cannot find neither xxd nor python" "$LINENO" 5
+    fi
+    ;;
+esac
+
 # Save it here so that, even in case of --enable-libgcj, if the Java
 # front-end isn't enabled, we still get libgcj disabled.
 libgcj_saved=$libgcj
diff --git a/configure.ac b/configure.ac
index fbc49ce..f429a1f 100644
--- a/configure.ac
+++ b/configure.ac
@@ -494,6 +494,18 @@ else
 fi])
 AC_SUBST(extra_liboffloadmic_configure_flags)
 
+# Check if xxd is present in the system
+# when the target is intelmic or intelmicemul.
+case "${target}" in
+  *-intelmic-* | *-intelmicemul-*)
+    AC_CHECK_PROG(xxd_present, xxd, "yes", "no")
+    AC_CHECK_PROG(python_present, python, "yes", "no")
+    if test "$xxd_present$python_present" = "nono"; then
+      AC_MSG_ERROR([cannot find neither xxd nor python])
+    fi
+    ;;
+esac
+
 # Save it here so that, even in case of --enable-libgcj, if the Java
 # front-end isn't enabled, we still get libgcj disabled.
 libgcj_saved=$libgcj
diff --git a/liboffloadmic/plugin/Makefile.am b/liboffloadmic/plugin/Makefile.am
index 19d69ab..cb93309 100644
--- a/liboffloadmic/plugin/Makefile.am
+++ b/liboffloadmic/plugin/Makefile.am
@@ -49,6 +49,10 @@ target_prefix_dir = $(libdir)/gcc/$(accel_target)/$(gcc_version)$(MULTISUBDIR)
 target_build_dir = $(accel_search_dir)/$(accel_target)$(MULTISUBDIR)/liboffloadmic/plugin
 target_install_dir = $(accel_search_dir)/lib/gcc/$(accel_target)/$(gcc_version)$(MULTISUBDIR)
 
+XXD_PY = $(source_dir)/xxd.py
+xxd_path=`which xxd`
+python_path=`which python`
+
 if PLUGIN_HOST
   toolexeclib_LTLIBRARIES = libgomp-plugin-intelmic.la
   libgomp_plugin_intelmic_la_SOURCES = libgomp-plugin-intelmic.cpp
@@ -73,7 +77,7 @@ main_target_image.h: offload_target_main
 	@echo "};" >> $@
 	@echo "extern \"C\" const MainTargetImage main_target_image = {" >> $@
 	@echo "  image_size, \"offload_target_main\"," >> $@
-	@cat $< | xxd -include >> $@
+	@if test "x$(xxd_path)" != "x"; then cat $< | $(xxd_path) -include >> $@; else $(python_path) $(XXD_PY) $< >> $@; fi;
 	@echo "};" >> $@
 
 offload_target_main: $(liboffload_dir)/ofldbegin.o offload_target_main.o $(liboffload_dir)/ofldend.o
diff --git a/liboffloadmic/plugin/Makefile.in b/liboffloadmic/plugin/Makefile.in
index 19a1a96..12fe5ca 100644
--- a/liboffloadmic/plugin/Makefile.in
+++ b/liboffloadmic/plugin/Makefile.in
@@ -318,6 +318,9 @@ libsubincludedir = $(libdir)/gcc/$(target_alias)/$(gcc_version)$(MULTISUBDIR)/in
 target_prefix_dir = $(libdir)/gcc/$(accel_target)/$(gcc_version)$(MULTISUBDIR)
 target_build_dir = $(accel_search_dir)/$(accel_target)$(MULTISUBDIR)/liboffloadmic/plugin
 target_install_dir = $(accel_search_dir)/lib/gcc/$(accel_target)/$(gcc_version)$(MULTISUBDIR)
+XXD_PY = $(source_dir)/xxd.py
+xxd_path = `which xxd`
+python_path = `which python`
 @PLUGIN_HOST_TRUE@toolexeclib_LTLIBRARIES = libgomp-plugin-intelmic.la
 @PLUGIN_HOST_TRUE@libgomp_plugin_intelmic_la_SOURCES = libgomp-plugin-intelmic.cpp
 @PLUGIN_HOST_TRUE@libgomp_plugin_intelmic_la_CPPFLAGS = $(CPPFLAGS) -DLINUX -DCOI_LIBRARY_VERSION=2 -DMYO_SUPPORT -DOFFLOAD_DEBUG=1 -DSEP_SUPPORT -DTIMING_SUPPORT -DHOST_LIBRARY=1 -I$(coi_inc_dir) -I$(myo_inc_dir) -I$(liboffload_src_dir) -I$(libgomp_src_dir) -I$(libgomp_dir) -I$(target_prefix_dir)/include -I$(target_build_dir) -I$(target_install_dir)/include
@@ -719,7 +722,7 @@ main_target_image.h: offload_target_main
 	@echo "};" >> $@
 	@echo "extern \"C\" const MainTargetImage main_target_image = {" >> $@
 	@echo "  image_size, \"offload_target_main\"," >> $@
-	@cat $< | xxd -include >> $@
+	@if test "x$(xxd_path)" != "x"; then cat $< | $(xxd_path) -include >> $@; else $(python_path) $(XXD_PY) $< >> $@; fi;
 	@echo "};" >> $@
 
 offload_target_main: $(liboffload_dir)/ofldbegin.o offload_target_main.o $(liboffload_dir)/ofldend.o
diff --git a/liboffloadmic/plugin/xxd.py b/liboffloadmic/plugin/xxd.py
new file mode 100644
index 0000000..fd9d7f0
--- /dev/null
+++ b/liboffloadmic/plugin/xxd.py
@@ -0,0 +1,39 @@
+#!/usr/bin/python
+import sys
+
+if sys.version_info[0] == 2:
+    # Python 2:
+    # 'block' below is an instance of str; iterating over it gives us
+    # str instances of len 1.
+    def get_byte(ch):
+        return ord(ch)
+else:
+    # Python 3:
+    # 'block' below is an instance of bytes; iterating over it gives us
+    # instances of int, in the range 0-255.
+    def get_byte(ch):
+        return ch
+
+with open(sys.argv[1],"rb") as f:
+    nextblock = f.read(12)
+    while 1:
+        block = nextblock
+        nextblock = f.read(12)
+        if not block:
+            break
+        str = ""
+        for item in block:
+            byte = get_byte(item)
+            if str == "":
+                str = "  "
+            else:
+                str += ", "
+            if byte < 10:
+                str += "0x0" + chr(ord('0')+byte)
+            elif byte < 16:
+                str += "0x0" + chr(ord('a')+byte-10)
+            else:
+                str += hex(byte)
+        if nextblock:
+            str += ","
+        print(str)

Re: [PATCH 3/4] Add libgomp plugin for Intel MIC

[Maxim Blumental]

Review the patches in the previous letter, please.

2015-08-06 17:34 GMT+03:00 Maxim Blumental <bvmaks@gmail.com>:
>  Applied the idea with python script alternative. Review, please.



-- 


---------------------
Sincerely yours,
Maxim Blumental