* Re: [PATCH 3/4] Add libgomp plugin for Intel MIC [not found] <20151111145615.GA4807@msticlxl57.ims.intel.com> @ 2015-11-30 15:11 ` Aleksander Ivanyushenko 2015-11-30 15:37 ` Jakub Jelinek 1 sibling, 0 replies; 24+ messages in thread From: Aleksander Ivanyushenko @ 2015-11-30 15:11 UTC (permalink / raw) To: jakub, kirill.yukhin, iverbin, gcc-patches 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. ^ permalink raw reply [flat|nested] 24+ messages in thread
* Re: [PATCH 3/4] Add libgomp plugin for Intel MIC [not found] <20151111145615.GA4807@msticlxl57.ims.intel.com> 2015-11-30 15:11 ` [PATCH 3/4] Add libgomp plugin for Intel MIC Aleksander Ivanyushenko @ 2015-11-30 15:37 ` Jakub Jelinek 1 sibling, 0 replies; 24+ messages in thread From: Jakub Jelinek @ 2015-11-30 15:37 UTC (permalink / raw) To: Aleksander Ivanyushenko; +Cc: kirill.yukhin, iverbin, gcc-patches 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 ^ permalink raw reply [flat|nested] 24+ messages in thread
* [PATCH 0/4] OpenMP 4.0 offloading to Intel MIC @ 2014-10-21 17:16 Ilya Verbin 2014-10-21 17:28 ` [PATCH 3/4] Add libgomp plugin for " Ilya Verbin 0 siblings, 1 reply; 24+ messages in thread From: Ilya Verbin @ 2014-10-21 17:16 UTC (permalink / raw) To: gcc-patches 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 ^ permalink raw reply [flat|nested] 24+ messages in thread
* [PATCH 3/4] Add libgomp plugin for Intel MIC 2014-10-21 17:16 [PATCH 0/4] OpenMP 4.0 offloading to " Ilya Verbin @ 2014-10-21 17:28 ` Ilya Verbin 2014-10-22 9:47 ` Jakub Jelinek 2015-07-08 14:16 ` Thomas Schwinge 0 siblings, 2 replies; 24+ messages in thread From: Ilya Verbin @ 2014-10-21 17:28 UTC (permalink / raw) To: Jakub Jelinek, gcc-patches; +Cc: Kirill Yukhin, Andrey Turetskiy 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 ^ permalink raw reply [flat|nested] 24+ messages in thread
* Re: [PATCH 3/4] Add libgomp plugin for Intel MIC 2014-10-21 17:28 ` [PATCH 3/4] Add libgomp plugin for " Ilya Verbin @ 2014-10-22 9:47 ` Jakub Jelinek 2014-10-23 16:00 ` Ilya Verbin 2015-07-08 14:16 ` Thomas Schwinge 1 sibling, 1 reply; 24+ messages in thread From: Jakub Jelinek @ 2014-10-22 9:47 UTC (permalink / raw) To: Ilya Verbin; +Cc: gcc-patches, Kirill Yukhin, Andrey Turetskiy 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 ^ permalink raw reply [flat|nested] 24+ messages in thread
* Re: [PATCH 3/4] Add libgomp plugin for Intel MIC 2014-10-22 9:47 ` Jakub Jelinek @ 2014-10-23 16:00 ` Ilya Verbin 2014-10-24 14:57 ` Jakub Jelinek 0 siblings, 1 reply; 24+ messages in thread From: Ilya Verbin @ 2014-10-23 16:00 UTC (permalink / raw) To: Jakub Jelinek; +Cc: gcc-patches, Kirill Yukhin, Andrey Turetskiy 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 ^ permalink raw reply [flat|nested] 24+ messages in thread
* Re: [PATCH 3/4] Add libgomp plugin for Intel MIC 2014-10-23 16:00 ` Ilya Verbin @ 2014-10-24 14:57 ` Jakub Jelinek 2014-10-24 15:12 ` Ilya Verbin 0 siblings, 1 reply; 24+ messages in thread From: Jakub Jelinek @ 2014-10-24 14:57 UTC (permalink / raw) To: Ilya Verbin; +Cc: gcc-patches, Kirill Yukhin, Andrey Turetskiy 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 ^ permalink raw reply [flat|nested] 24+ messages in thread
* Re: [PATCH 3/4] Add libgomp plugin for Intel MIC 2014-10-24 14:57 ` Jakub Jelinek @ 2014-10-24 15:12 ` Ilya Verbin 2014-10-24 15:19 ` Jakub Jelinek 0 siblings, 1 reply; 24+ messages in thread From: Ilya Verbin @ 2014-10-24 15:12 UTC (permalink / raw) To: Jakub Jelinek; +Cc: gcc-patches, Kirill Yukhin, Andrey Turetskiy 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 ^ permalink raw reply [flat|nested] 24+ messages in thread
* Re: [PATCH 3/4] Add libgomp plugin for Intel MIC 2014-10-24 15:12 ` Ilya Verbin @ 2014-10-24 15:19 ` Jakub Jelinek 2014-10-27 14:24 ` Ilya Verbin 0 siblings, 1 reply; 24+ messages in thread From: Jakub Jelinek @ 2014-10-24 15:19 UTC (permalink / raw) To: Ilya Verbin; +Cc: gcc-patches, Kirill Yukhin, Andrey Turetskiy 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 ^ permalink raw reply [flat|nested] 24+ messages in thread
* Re: [PATCH 3/4] Add libgomp plugin for Intel MIC 2014-10-24 15:19 ` Jakub Jelinek @ 2014-10-27 14:24 ` Ilya Verbin 2014-11-06 18:25 ` Jakub Jelinek 0 siblings, 1 reply; 24+ messages in thread From: Ilya Verbin @ 2014-10-27 14:24 UTC (permalink / raw) To: Jakub Jelinek; +Cc: gcc-patches, Kirill Yukhin, Andrey Turetskiy 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 ^ permalink raw reply [flat|nested] 24+ messages in thread
* Re: [PATCH 3/4] Add libgomp plugin for Intel MIC 2014-10-27 14:24 ` Ilya Verbin @ 2014-11-06 18:25 ` Jakub Jelinek 2014-11-10 14:32 ` Ilya Verbin 0 siblings, 1 reply; 24+ messages in thread From: Jakub Jelinek @ 2014-11-06 18:25 UTC (permalink / raw) To: Ilya Verbin; +Cc: gcc-patches, Kirill Yukhin, Andrey Turetskiy 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 ^ permalink raw reply [flat|nested] 24+ messages in thread
* Re: [PATCH 3/4] Add libgomp plugin for Intel MIC 2014-11-06 18:25 ` Jakub Jelinek @ 2014-11-10 14:32 ` Ilya Verbin 2014-11-11 7:07 ` Jakub Jelinek 2014-12-12 9:42 ` Thomas Schwinge 0 siblings, 2 replies; 24+ messages in thread From: Ilya Verbin @ 2014-11-10 14:32 UTC (permalink / raw) To: Jakub Jelinek; +Cc: gcc-patches, Kirill Yukhin, Andrey Turetskiy 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 ^ permalink raw reply [flat|nested] 24+ messages in thread
* Re: [PATCH 3/4] Add libgomp plugin for Intel MIC 2014-11-10 14:32 ` Ilya Verbin @ 2014-11-11 7:07 ` Jakub Jelinek 2014-12-12 9:42 ` Thomas Schwinge 1 sibling, 0 replies; 24+ messages in thread From: Jakub Jelinek @ 2014-11-11 7:07 UTC (permalink / raw) To: Ilya Verbin; +Cc: gcc-patches, Kirill Yukhin, Andrey Turetskiy 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 ^ permalink raw reply [flat|nested] 24+ messages in thread
* Re: [PATCH 3/4] Add libgomp plugin for Intel MIC 2014-11-10 14:32 ` Ilya Verbin 2014-11-11 7:07 ` Jakub Jelinek @ 2014-12-12 9:42 ` Thomas Schwinge 2015-01-08 14:48 ` Thomas Schwinge 1 sibling, 1 reply; 24+ messages in thread From: Thomas Schwinge @ 2014-12-12 9:42 UTC (permalink / raw) To: Ilya Verbin, Jakub Jelinek; +Cc: gcc-patches, Kirill Yukhin, Andrey Turetskiy [-- 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 --] ^ permalink raw reply [flat|nested] 24+ messages in thread
* Re: [PATCH 3/4] Add libgomp plugin for Intel MIC 2014-12-12 9:42 ` Thomas Schwinge @ 2015-01-08 14:48 ` Thomas Schwinge 0 siblings, 0 replies; 24+ messages in thread From: Thomas Schwinge @ 2015-01-08 14:48 UTC (permalink / raw) To: Ilya Verbin, Jakub Jelinek; +Cc: gcc-patches, Kirill Yukhin, Andrey Turetskiy [-- 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 --] ^ permalink raw reply [flat|nested] 24+ messages in thread
* Re: [PATCH 3/4] Add libgomp plugin for Intel MIC 2014-10-21 17:28 ` [PATCH 3/4] Add libgomp plugin for " Ilya Verbin 2014-10-22 9:47 ` Jakub Jelinek @ 2015-07-08 14:16 ` Thomas Schwinge 2015-07-08 15:14 ` Ilya Verbin 2015-07-23 19:05 ` Ilya Verbin 1 sibling, 2 replies; 24+ messages in thread From: Thomas Schwinge @ 2015-07-08 14:16 UTC (permalink / raw) To: Ilya Verbin, Jakub Jelinek, gcc-patches; +Cc: Kirill Yukhin, Andrey Turetskiy [-- 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 --] ^ permalink raw reply [flat|nested] 24+ messages in thread
* Re: [PATCH 3/4] Add libgomp plugin for Intel MIC 2015-07-08 14:16 ` Thomas Schwinge @ 2015-07-08 15:14 ` Ilya Verbin 2015-07-08 15:52 ` Thomas Schwinge 2015-07-23 19:05 ` Ilya Verbin 1 sibling, 1 reply; 24+ messages in thread From: Ilya Verbin @ 2015-07-08 15:14 UTC (permalink / raw) To: Thomas Schwinge Cc: Jakub Jelinek, <gcc-patches@gcc.gnu.org>, Kirill Yukhin, Andrey Turetskiy > 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 ^ permalink raw reply [flat|nested] 24+ messages in thread
* Re: [PATCH 3/4] Add libgomp plugin for Intel MIC 2015-07-08 15:14 ` Ilya Verbin @ 2015-07-08 15:52 ` Thomas Schwinge 0 siblings, 0 replies; 24+ messages in thread From: Thomas Schwinge @ 2015-07-08 15:52 UTC (permalink / raw) To: Ilya Verbin Cc: Jakub Jelinek, <gcc-patches@gcc.gnu.org>, Kirill Yukhin, Andrey Turetskiy [-- 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 --] ^ permalink raw reply [flat|nested] 24+ messages in thread
* Re: [PATCH 3/4] Add libgomp plugin for Intel MIC 2015-07-08 14:16 ` Thomas Schwinge 2015-07-08 15:14 ` Ilya Verbin @ 2015-07-23 19:05 ` Ilya Verbin 2015-07-24 8:06 ` Jakub Jelinek 1 sibling, 1 reply; 24+ messages in thread From: Ilya Verbin @ 2015-07-23 19:05 UTC (permalink / raw) To: Thomas Schwinge, Jakub Jelinek; +Cc: gcc-patches, Kirill Yukhin, bvmaks 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 ^ permalink raw reply [flat|nested] 24+ messages in thread
* Re: [PATCH 3/4] Add libgomp plugin for Intel MIC 2015-07-23 19:05 ` Ilya Verbin @ 2015-07-24 8:06 ` Jakub Jelinek 2015-07-24 14:27 ` David Malcolm 0 siblings, 1 reply; 24+ messages in thread From: Jakub Jelinek @ 2015-07-24 8:06 UTC (permalink / raw) To: Ilya Verbin, David Malcolm Cc: Thomas Schwinge, gcc-patches, Kirill Yukhin, bvmaks 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 ^ permalink raw reply [flat|nested] 24+ messages in thread
* Re: [PATCH 3/4] Add libgomp plugin for Intel MIC 2015-07-24 8:06 ` Jakub Jelinek @ 2015-07-24 14:27 ` David Malcolm 2015-07-28 15:51 ` Maxim Blumental 0 siblings, 1 reply; 24+ messages in thread From: David Malcolm @ 2015-07-24 14:27 UTC (permalink / raw) To: Jakub Jelinek Cc: Ilya Verbin, Thomas Schwinge, gcc-patches, Kirill Yukhin, bvmaks [-- 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) ^ permalink raw reply [flat|nested] 24+ messages in thread
* Re: [PATCH 3/4] Add libgomp plugin for Intel MIC 2015-07-24 14:27 ` David Malcolm @ 2015-07-28 15:51 ` Maxim Blumental 2015-08-03 10:24 ` Maxim Blumental 2015-08-06 14:35 ` Fwd: " Maxim Blumental 0 siblings, 2 replies; 24+ messages in thread From: Maxim Blumental @ 2015-07-28 15:51 UTC (permalink / raw) To: David Malcolm Cc: Jakub Jelinek, Ilya Verbin, Thomas Schwinge, gcc-patches, Kirill Yukhin [-- 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) ^ permalink raw reply [flat|nested] 24+ messages in thread
* Re: [PATCH 3/4] Add libgomp plugin for Intel MIC 2015-07-28 15:51 ` Maxim Blumental @ 2015-08-03 10:24 ` Maxim Blumental 2015-08-04 17:40 ` David Malcolm 2015-08-06 14:35 ` Fwd: " Maxim Blumental 1 sibling, 1 reply; 24+ messages in thread From: Maxim Blumental @ 2015-08-03 10:24 UTC (permalink / raw) To: David Malcolm Cc: Jakub Jelinek, Ilya Verbin, Thomas Schwinge, gcc-patches, Kirill Yukhin 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 ^ permalink raw reply [flat|nested] 24+ messages in thread
* Re: [PATCH 3/4] Add libgomp plugin for Intel MIC 2015-08-03 10:24 ` Maxim Blumental @ 2015-08-04 17:40 ` David Malcolm 0 siblings, 0 replies; 24+ messages in thread From: David Malcolm @ 2015-08-04 17:40 UTC (permalink / raw) To: Maxim Blumental Cc: Jakub Jelinek, Ilya Verbin, Thomas Schwinge, gcc-patches, Kirill Yukhin 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 ^ permalink raw reply [flat|nested] 24+ messages in thread
* Fwd: [PATCH 3/4] Add libgomp plugin for Intel MIC 2015-07-28 15:51 ` Maxim Blumental 2015-08-03 10:24 ` Maxim Blumental @ 2015-08-06 14:35 ` Maxim Blumental 2015-08-11 12:27 ` Maxim Blumental 1 sibling, 1 reply; 24+ messages in thread From: Maxim Blumental @ 2015-08-06 14:35 UTC (permalink / raw) To: Jakub Jelinek; +Cc: David Malcolm, Kirill Yukhin, Ilya Verbin, gcc-patches [-- 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) ^ permalink raw reply [flat|nested] 24+ messages in thread
* Re: [PATCH 3/4] Add libgomp plugin for Intel MIC 2015-08-06 14:35 ` Fwd: " Maxim Blumental @ 2015-08-11 12:27 ` Maxim Blumental 0 siblings, 0 replies; 24+ messages in thread From: Maxim Blumental @ 2015-08-11 12:27 UTC (permalink / raw) To: Jakub Jelinek; +Cc: David Malcolm, Kirill Yukhin, Ilya Verbin, gcc-patches 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 ^ permalink raw reply [flat|nested] 24+ messages in thread
end of thread, other threads:[~2015-11-30 15:23 UTC | newest] Thread overview: 24+ messages (download: mbox.gz / follow: Atom feed) -- links below jump to the message on this page -- [not found] <20151111145615.GA4807@msticlxl57.ims.intel.com> 2015-11-30 15:11 ` [PATCH 3/4] Add libgomp plugin for Intel MIC Aleksander Ivanyushenko 2015-11-30 15:37 ` Jakub Jelinek 2014-10-21 17:16 [PATCH 0/4] OpenMP 4.0 offloading to " Ilya Verbin 2014-10-21 17:28 ` [PATCH 3/4] Add libgomp plugin for " Ilya Verbin 2014-10-22 9:47 ` Jakub Jelinek 2014-10-23 16:00 ` Ilya Verbin 2014-10-24 14:57 ` Jakub Jelinek 2014-10-24 15:12 ` Ilya Verbin 2014-10-24 15:19 ` Jakub Jelinek 2014-10-27 14:24 ` Ilya Verbin 2014-11-06 18:25 ` Jakub Jelinek 2014-11-10 14:32 ` Ilya Verbin 2014-11-11 7:07 ` Jakub Jelinek 2014-12-12 9:42 ` Thomas Schwinge 2015-01-08 14:48 ` Thomas Schwinge 2015-07-08 14:16 ` Thomas Schwinge 2015-07-08 15:14 ` Ilya Verbin 2015-07-08 15:52 ` Thomas Schwinge 2015-07-23 19:05 ` Ilya Verbin 2015-07-24 8:06 ` Jakub Jelinek 2015-07-24 14:27 ` David Malcolm 2015-07-28 15:51 ` Maxim Blumental 2015-08-03 10:24 ` Maxim Blumental 2015-08-04 17:40 ` David Malcolm 2015-08-06 14:35 ` Fwd: " Maxim Blumental 2015-08-11 12:27 ` Maxim Blumental
This is a public inbox, see mirroring instructions for how to clone and mirror all data and code used for this inbox; as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).