From patchwork Thu Sep 19 16:09:04 2013 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Michael Zolotukhin X-Patchwork-Id: 275998 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Received: from sourceware.org (server1.sourceware.org [209.132.180.131]) (using TLSv1 with cipher DHE-RSA-AES256-SHA (256/256 bits)) (Client did not present a certificate) by ozlabs.org (Postfix) with ESMTPS id 2CD4E2C0110 for ; Fri, 20 Sep 2013 02:10:17 +1000 (EST) DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:date :from:to:cc:subject:message-id:references:mime-version :content-type:in-reply-to; q=dns; s=default; b=YZPPsYPB4a2DlTZX3 uLhmceGFfmsawypIG5YiSY5BBnfcc78u2xlPvyQQ0RduGtsYb9fzT8ZLTh6nqaUv J8ui+hfa4/gm6nKA8utuw0nl6o6H56awdbgRADEwc2ntXHAzqlt19bdYMdJ8j0hF JtfqlTm6o6wz7G2oYTdgK+4CIE= DKIM-Signature: v=1; a=rsa-sha1; c=relaxed; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:date :from:to:cc:subject:message-id:references:mime-version :content-type:in-reply-to; s=default; bh=EhA1Y3KC80TpswNR669/sKx fVYQ=; b=TzLdJ9z1lB+nL6VdVvXHjD8LjWUIS10rVQH2W4Ga3Uxsvup+GgSNJr8 eQYzzcTji3tyWfEHHyLgv4QCbJlYyzakaeCumkIbi9LABdUmxL5k0/kEhS/cycqR GAtviwwf37Cb8gt3CCuylphc/nR+vEQAkAD87DgW25JyL7F9Wyj8= Received: (qmail 24760 invoked by alias); 19 Sep 2013 16:10:09 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Unsubscribe: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Delivered-To: mailing list gcc-patches@gcc.gnu.org Received: (qmail 24746 invoked by uid 89); 19 Sep 2013 16:10:08 -0000 Received: from mail-pd0-f170.google.com (HELO mail-pd0-f170.google.com) (209.85.192.170) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with (AES128-SHA encrypted) ESMTPS; Thu, 19 Sep 2013 16:10:08 +0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-2.9 required=5.0 tests=ALL_TRUSTED, AWL, BAYES_00, FREEMAIL_FROM autolearn=ham version=3.3.2 X-HELO: mail-pd0-f170.google.com Received: by mail-pd0-f170.google.com with SMTP id x10so8665286pdj.1 for ; Thu, 19 Sep 2013 09:09:35 -0700 (PDT) X-Received: by 10.68.213.167 with SMTP id nt7mr2815102pbc.140.1379606975738; Thu, 19 Sep 2013 09:09:35 -0700 (PDT) Received: from msticlxl57.ims.intel.com ([192.55.54.42]) by mx.google.com with ESMTPSA id bt1sm10217443pbb.2.1969.12.31.16.00.00 (version=TLSv1 cipher=RC4-SHA bits=128/128); Thu, 19 Sep 2013 09:09:34 -0700 (PDT) Date: Thu, 19 Sep 2013 20:09:04 +0400 From: "Michael V. Zolotukhin" To: Jakub Jelinek Cc: Richard Henderson , Kirill Yukhin , gcc Subject: Re: [PATCH][gomp4] Plugins Support in LibGOMP (Take 2) Message-ID: <20130919160904.GD11801@msticlxl57.ims.intel.com> References: <20130918083538.GA11801@msticlxl57.ims.intel.com> <20130918090525.GF1817@tucnak.redhat.com> MIME-Version: 1.0 Content-Disposition: inline In-Reply-To: <20130918090525.GF1817@tucnak.redhat.com> User-Agent: Mutt/1.5.21 (2010-09-15) X-IsSubscribed: yes Hi Jakub, Updated patch and my answers are below. > The OpenMP standard has the omp_is_initial_device () function that can be > used to query whether the code is offloaded or not. So I don't think we > need to do the logging. For the device 257 hack we of course don't return > that as true, but that is a hack that is going away. Ok that sounds good too. > > @@ -50,6 +59,10 @@ struct target_mem_desc { > > struct target_mem_desc *prev; > > /* Number of items in following list. */ > > size_t list_count; > > + > > + /* Corresponding target device descriptor. */ > > + struct gomp_device_descr* device_descr; > > Please put the space before *, not after it. I wasn't aware of that rule, sorry. Fixed. > > + /* Plugin file name. */ > > + char plugin_name[PATH_MAX]; > > I don't like such fixed size arrays, for most cases > it will be big memory waste. What do you need the plugin_name > for? And, if you really need it past dlopen, can't you store > it as const char *plugin_name instead? I kept it just in case - it easily could be removed, and I did it in the current version of the patch. > > + > > + /* Plugin file handler. */ > > + void *plugin_handle; > > + > > + /* Function handlers. */ > > + bool (*device_available_func) (void); > > The scan hook shouldn't give you just bool whether the device is available, > but how many devices of that kind are available. You can have 2 MIC > cards and one or two HSAIL GPGPU in a box e.g. Plus, is this hook useful > after the initialization at all? I'd say it would be enough to just > dlsym it during initialization, ask how many devices it has and just create > that many device structures with that plugin_handle. > What you want are hooks for device_alloc (taking size and align arguments, > returning uintptr_t target address), device_free (taking uintptr_t target > address and perhaps size), device_copyto (like memcpy, just with target > address uintptr_t instead of void *) and device_copyfrom (similarly), > and device_run hook or similar (taking host and target fn and target > uintptr_t address of the block with pointers). That's just a stub, showing how everything would work in future, when the interface libgomp<->plugin would be finally settled. I think it's better to wait a little bit when we would progress further in development of the libgomp plugin - probably we'd spot new issues in the interface. Anyway, it's easy to add any routines we want here. > You need to call pthread_once here too, so that omp_get_num_devices returns > the correct number. > ... > Thus, IMHO you should just call gomp_get_num_devices () here, or after the > if (device_id == -1) block, and that will ensure gomp_target_init has been > already called. Just save the return value into a temporary. Fixed. > > > + if (device_id == -1) > > { > > struct gomp_task_icv *icv = gomp_icv (false); > > - device = icv->default_device_var; > > + device_id = icv->default_device_var; > > } > > /* FIXME: Temporary hack for testing non-shared address spaces on host. */ > > - if (device == 257) > > - return 257; > > - if (device >= gomp_get_num_devices ()) > > - return -1; > > - return -1; > > + if (device_id == 257) > > + return &devices[0]; > > Guess the hack should be if gomp_get_num_devices () returned 0 and > device_id == 257, otherwise the hack device won't be created. Currently we always have at least one device (see FIXME in gomp_find_available_plugins routine) - even if we found no plugins, we create a hack device. If we found some plugins, then we don't create a new device for the hack, but use the devices[0] for it. > > - struct target_mem_desc *tgt > > - = gomp_malloc (sizeof (*tgt) + sizeof (tgt->list[0]) * mapnum); > > + struct target_mem_desc *tgt = NULL; > > + tgt = gomp_malloc (sizeof (*tgt) + sizeof (tgt->list[0]) * mapnum); > > Why this change? Changed back. > > tgt->list_count = mapnum; > > tgt->refcount = 1; > > + tgt->device_descr = devicep; > > + > > + if (!devicep) > > + return tgt; > > Why this conditional? mapnum == 0 conditional below will do the trick. Fixed. > > + /* FIXME: currently only device 257 is available and it is a hack which is > > + done only to test the functionality early. We need to enable all devices, > > + not only this one. */ > > Yeah, I don't see why the FIXME is here, just use gomp_map_vars > unconditionally, or conditionally on some flag in the device descr structure > (whether device has non-shared address space). Removed. > > + if (devicep->id == 257) > > { > > struct target_mem_desc *tgt > > - = gomp_map_vars (mapnum, hostaddrs, sizes, kinds, true); > > + = gomp_map_vars (devicep, mapnum, hostaddrs, sizes, kinds, true); > > fn ((void *) tgt->tgt_start); > > And thus would be devicep->device_run hook. We'll start device_run hook here once the interface libgomp<->plugin is fully set. > Why devicep here, when you know it is NULL? Fixed. > > + strncpy (current_device.plugin_name, plugin_path, PATH_MAX); > > + strcat (current_device.plugin_name, "/"); > > + strcat (current_device.plugin_name, ent->d_name); > > Potential buffer overflow. Fixed. > > +/* This function initializes runtime needed for offloading. > > + It loads plugins, sets up a connection with devices, etc. */ > > +static void > > +gomp_target_init (void) > > +{ > > + gomp_find_available_plugins (); > > +} > > Why this indirection? Just rename gomp_find_available_plugins to > gomp_target_init? I think we might want to do something else in gomp_target_init, not just look for available plugins. If it's not changed in future, then yes, we could just rename gomp_find_available_plugins to gomp_target_init. Michael > Jakub --- libgomp/config.h.in | 6 + libgomp/configure | 63 +++++++++++ libgomp/configure.ac | 9 ++ libgomp/target.c | 303 +++++++++++++++++++++++++++++++++++++++++---------- 4 files changed, 324 insertions(+), 57 deletions(-) diff --git a/libgomp/config.h.in b/libgomp/config.h.in index 14c7e2a..67f5420 100644 --- a/libgomp/config.h.in +++ b/libgomp/config.h.in @@ -30,6 +30,9 @@ /* Define to 1 if you have the header file. */ #undef HAVE_INTTYPES_H +/* Define to 1 if you have the `dl' library (-ldl). */ +#undef HAVE_LIBDL + /* Define to 1 if you have the header file. */ #undef HAVE_MEMORY_H @@ -107,6 +110,9 @@ /* Define to the version of this package. */ #undef PACKAGE_VERSION +/* Define if all infrastructure, needed for plugins, is supported. */ +#undef PLUGIN_SUPPORT + /* The size of `char', as computed by sizeof. */ #undef SIZEOF_CHAR diff --git a/libgomp/configure b/libgomp/configure index 238b1af..f4f71a4 100755 --- a/libgomp/configure +++ b/libgomp/configure @@ -15046,6 +15046,69 @@ fi rm -f core conftest.err conftest.$ac_objext \ conftest$ac_exeext conftest.$ac_ext +plugin_support=yes +{ $as_echo "$as_me:${as_lineno-$LINENO}: checking for dlsym in -ldl" >&5 +$as_echo_n "checking for dlsym in -ldl... " >&6; } +if test "${ac_cv_lib_dl_dlsym+set}" = set; then : + $as_echo_n "(cached) " >&6 +else + ac_check_lib_save_LIBS=$LIBS +LIBS="-ldl $LIBS" +cat confdefs.h - <<_ACEOF >conftest.$ac_ext +/* end confdefs.h. */ + +/* Override any GCC internal prototype to avoid an error. + Use char because int might match the return type of a GCC + builtin and then its argument prototype would still apply. */ +#ifdef __cplusplus +extern "C" +#endif +char dlsym (); +int +main () +{ +return dlsym (); + ; + return 0; +} +_ACEOF +if ac_fn_c_try_link "$LINENO"; then : + ac_cv_lib_dl_dlsym=yes +else + ac_cv_lib_dl_dlsym=no +fi +rm -f core conftest.err conftest.$ac_objext \ + conftest$ac_exeext conftest.$ac_ext +LIBS=$ac_check_lib_save_LIBS +fi +{ $as_echo "$as_me:${as_lineno-$LINENO}: result: $ac_cv_lib_dl_dlsym" >&5 +$as_echo "$ac_cv_lib_dl_dlsym" >&6; } +if test "x$ac_cv_lib_dl_dlsym" = x""yes; then : + cat >>confdefs.h <<_ACEOF +#define HAVE_LIBDL 1 +_ACEOF + + LIBS="-ldl $LIBS" + +else + plugin_support=no +fi + +ac_fn_c_check_header_mongrel "$LINENO" "dirent.h" "ac_cv_header_dirent_h" "$ac_includes_default" +if test "x$ac_cv_header_dirent_h" = x""yes; then : + +else + plugin_support=no +fi + + + +if test x$plugin_support = xyes; then + +$as_echo "#define PLUGIN_SUPPORT 1" >>confdefs.h + +fi + # Check for functions needed. for ac_func in getloadavg clock_gettime strtoull do : diff --git a/libgomp/configure.ac b/libgomp/configure.ac index d87ed29..85ecbcf 100644 --- a/libgomp/configure.ac +++ b/libgomp/configure.ac @@ -193,6 +193,15 @@ AC_LINK_IFELSE( [], [AC_MSG_ERROR([Pthreads are required to build libgomp])])]) +plugin_support=yes +AC_CHECK_LIB(dl, dlsym, , [plugin_support=no]) +AC_CHECK_HEADER(dirent.h, , [plugin_support=no]) + +if test x$plugin_support = xyes; then + AC_DEFINE(PLUGIN_SUPPORT, 1, + [Define if all infrastructure, needed for plugins, is supported.]) +fi + # Check for functions needed. AC_CHECK_FUNCS(getloadavg clock_gettime strtoull) diff --git a/libgomp/target.c b/libgomp/target.c index 8b445bc..e1213b1 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -30,6 +30,15 @@ #include #include +#ifdef PLUGIN_SUPPORT +# include +# include +#endif + +static void gomp_target_init (void); + +static pthread_once_t gomp_is_initialized = PTHREAD_ONCE_INIT; + /* Forward declaration for a node in the tree. */ typedef struct splay_tree_node_s *splay_tree_node; typedef struct splay_tree_s *splay_tree; @@ -50,6 +59,10 @@ struct target_mem_desc { struct target_mem_desc *prev; /* Number of items in following list. */ size_t list_count; + + /* Corresponding target device descriptor. */ + struct gomp_device_descr *device_descr; + /* List of splay keys to remove (or decrease refcount) at the end of region. */ splay_tree_key list[]; @@ -70,6 +83,12 @@ struct splay_tree_key_s { bool copy_from; }; +/* Array of descriptors of all available devices. */ +static struct gomp_device_descr *devices; + +/* Total number of available devices. */ +static int num_devices; + /* The comparison function. */ static int @@ -87,33 +106,55 @@ splay_compare (splay_tree_key x, splay_tree_key y) #include "splay-tree.h" +/* This structure describes accelerator device. + It contains name of the corresponding libgomp plugin, function handlers for + interaction with the device, ID-number of the device, and information about + mapped memory. */ +struct gomp_device_descr +{ + /* This is the ID number of device. It could be specified in DEVICE-clause of + TARGET construct. */ + int id; + + /* Plugin file handler. */ + void *plugin_handle; + + /* Function handlers. */ + bool (*device_available_func) (void); + + /* Splay tree containing information about mapped memory regions. */ + struct splay_tree_s dev_splay_tree; + + /* Mutex for operating with the splay tree and other shared structures. */ + gomp_mutex_t dev_env_lock; +}; + attribute_hidden int gomp_get_num_devices (void) { - /* FIXME: Scan supported accelerators when called the first time. */ - return 0; + (void) pthread_once (&gomp_is_initialized, gomp_target_init); + return num_devices; } -static int -resolve_device (int device) +static struct gomp_device_descr * +resolve_device (int device_id) { - if (device == -1) + if (device_id == -1) { struct gomp_task_icv *icv = gomp_icv (false); - device = icv->default_device_var; + device_id = icv->default_device_var; } + if (device_id >= gomp_get_num_devices () + && device_id != 257) + return NULL; + /* FIXME: Temporary hack for testing non-shared address spaces on host. */ - if (device == 257) - return 257; - if (device >= gomp_get_num_devices ()) - return -1; - return -1; + if (device_id == 257) + return &devices[0]; + + return &devices[device_id]; } -/* These variables would be per-accelerator (which doesn't have shared address - space. */ -static struct splay_tree_s dev_splay_tree; -static gomp_mutex_t dev_env_lock; /* Handle the case where splay_tree_lookup found oldn for newn. Helper function of gomp_map_vars. */ @@ -137,8 +178,9 @@ gomp_map_vars_existing (splay_tree_key oldn, splay_tree_key newn, } static struct target_mem_desc * -gomp_map_vars (size_t mapnum, void **hostaddrs, size_t *sizes, - unsigned char *kinds, bool is_target) +gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum, + void **hostaddrs, size_t *sizes, unsigned char *kinds, + bool is_target) { size_t i, tgt_align, tgt_size, not_found_cnt = 0; struct splay_tree_key_s cur_node; @@ -146,6 +188,7 @@ gomp_map_vars (size_t mapnum, void **hostaddrs, size_t *sizes, = gomp_malloc (sizeof (*tgt) + sizeof (tgt->list[0]) * mapnum); tgt->list_count = mapnum; tgt->refcount = 1; + tgt->device_descr = devicep; if (mapnum == 0) return tgt; @@ -159,7 +202,7 @@ gomp_map_vars (size_t mapnum, void **hostaddrs, size_t *sizes, tgt_size = mapnum * sizeof (void *); } - gomp_mutex_lock (&dev_env_lock); + gomp_mutex_lock (&devicep->dev_env_lock); for (i = 0; i < mapnum; i++) { cur_node.host_start = (uintptr_t) hostaddrs[i]; @@ -167,7 +210,8 @@ gomp_map_vars (size_t mapnum, void **hostaddrs, size_t *sizes, cur_node.host_end = cur_node.host_start + sizes[i]; else cur_node.host_end = cur_node.host_start + sizeof (void *); - splay_tree_key n = splay_tree_lookup (&dev_splay_tree, &cur_node); + splay_tree_key n = splay_tree_lookup (&devicep->dev_splay_tree, + &cur_node); if (n) { tgt->list[i] = n; @@ -215,7 +259,7 @@ gomp_map_vars (size_t mapnum, void **hostaddrs, size_t *sizes, else k->host_end = k->host_start + sizeof (void *); splay_tree_key n - = splay_tree_lookup (&dev_splay_tree, k); + = splay_tree_lookup (&devicep->dev_splay_tree, k); if (n) { tgt->list[i] = n; @@ -235,7 +279,7 @@ gomp_map_vars (size_t mapnum, void **hostaddrs, size_t *sizes, tgt->refcount++; array->left = NULL; array->right = NULL; - splay_tree_insert (&dev_splay_tree, array); + splay_tree_insert (&devicep->dev_splay_tree, array); switch (kinds[i] & 7) { case 0: /* ALLOC */ @@ -257,16 +301,19 @@ gomp_map_vars (size_t mapnum, void **hostaddrs, size_t *sizes, /* Add bias to the pointer value. */ cur_node.host_start += sizes[i]; cur_node.host_end = cur_node.host_start + 1; - n = splay_tree_lookup (&dev_splay_tree, &cur_node); + n = splay_tree_lookup (&devicep->dev_splay_tree, + &cur_node); if (n == NULL) { /* Could be possibly zero size array section. */ cur_node.host_end--; - n = splay_tree_lookup (&dev_splay_tree, &cur_node); + n = splay_tree_lookup (&devicep->dev_splay_tree, + &cur_node); if (n == NULL) { cur_node.host_start--; - n = splay_tree_lookup (&dev_splay_tree, &cur_node); + n = splay_tree_lookup (&devicep->dev_splay_tree, + &cur_node); cur_node.host_start++; } } @@ -303,7 +350,7 @@ gomp_map_vars (size_t mapnum, void **hostaddrs, size_t *sizes, } } - gomp_mutex_unlock (&dev_env_lock); + gomp_mutex_unlock (&devicep->dev_env_lock); return tgt; } @@ -322,6 +369,8 @@ gomp_unmap_tgt (struct target_mem_desc *tgt) static void gomp_unmap_vars (struct target_mem_desc *tgt) { + struct gomp_device_descr *devicep = tgt->device_descr; + if (tgt->list_count == 0) { free (tgt); @@ -329,7 +378,7 @@ gomp_unmap_vars (struct target_mem_desc *tgt) } size_t i; - gomp_mutex_lock (&dev_env_lock); + gomp_mutex_lock (&devicep->dev_env_lock); for (i = 0; i < tgt->list_count; i++) if (tgt->list[i]->refcount > 1) tgt->list[i]->refcount--; @@ -341,7 +390,7 @@ gomp_unmap_vars (struct target_mem_desc *tgt) memcpy ((void *) k->host_start, (void *) (k->tgt->tgt_start + k->tgt_offset), k->host_end - k->host_start); - splay_tree_remove (&dev_splay_tree, k); + splay_tree_remove (&devicep->dev_splay_tree, k); if (k->tgt->refcount > 1) k->tgt->refcount--; else @@ -352,26 +401,30 @@ gomp_unmap_vars (struct target_mem_desc *tgt) tgt->refcount--; else gomp_unmap_tgt (tgt); - gomp_mutex_unlock (&dev_env_lock); + gomp_mutex_unlock (&devicep->dev_env_lock); } static void -gomp_update (size_t mapnum, void **hostaddrs, size_t *sizes, - unsigned char *kinds) +gomp_update (struct gomp_device_descr *devicep, size_t mapnum, + void **hostaddrs, size_t *sizes, unsigned char *kinds) { size_t i; struct splay_tree_key_s cur_node; + if (!devicep) + return; + if (mapnum == 0) return; - gomp_mutex_lock (&dev_env_lock); + gomp_mutex_lock (&devicep->dev_env_lock); for (i = 0; i < mapnum; i++) if (sizes[i]) { cur_node.host_start = (uintptr_t) hostaddrs[i]; cur_node.host_end = cur_node.host_start + sizes[i]; - splay_tree_key n = splay_tree_lookup (&dev_splay_tree, &cur_node); + splay_tree_key n = splay_tree_lookup (&devicep->dev_splay_tree, + &cur_node); if (n) { if (n->host_start > cur_node.host_start @@ -400,7 +453,7 @@ gomp_update (size_t mapnum, void **hostaddrs, size_t *sizes, (void *) cur_node.host_start, (void *) cur_node.host_end); } - gomp_mutex_unlock (&dev_env_lock); + gomp_mutex_unlock (&devicep->dev_env_lock); } /* Called when encountering a target directive. If DEVICE @@ -418,28 +471,26 @@ GOMP_target (int device, void (*fn) (void *), const void *openmp_target, size_t mapnum, void **hostaddrs, size_t *sizes, unsigned char *kinds) { - device = resolve_device (device); - if (device == -1) + struct gomp_device_descr *devicep = resolve_device (device); + if (devicep == NULL) { /* Host fallback. */ fn (hostaddrs); return; } - if (device == 257) - { - struct target_mem_desc *tgt - = gomp_map_vars (mapnum, hostaddrs, sizes, kinds, true); - fn ((void *) tgt->tgt_start); - gomp_unmap_vars (tgt); - } + + struct target_mem_desc *tgt + = gomp_map_vars (devicep, mapnum, hostaddrs, sizes, kinds, true); + fn ((void *) tgt->tgt_start); + gomp_unmap_vars (tgt); } void GOMP_target_data (int device, const void *openmp_target, size_t mapnum, void **hostaddrs, size_t *sizes, unsigned char *kinds) { - device = resolve_device (device); - if (device == -1) + struct gomp_device_descr *devicep = resolve_device (device); + if (devicep == NULL) { /* Host fallback. */ struct gomp_task_icv *icv = gomp_icv (false); @@ -450,21 +501,18 @@ GOMP_target_data (int device, const void *openmp_target, size_t mapnum, new #pragma omp target data, otherwise GOMP_target_end_data would get out of sync. */ struct target_mem_desc *tgt - = gomp_map_vars (0, NULL, NULL, NULL, false); + = gomp_map_vars (NULL, 0, NULL, NULL, NULL, false); tgt->prev = icv->target_data; icv->target_data = tgt; } return; } - if (device == 257) - { - struct target_mem_desc *tgt - = gomp_map_vars (mapnum, hostaddrs, sizes, kinds, false); - struct gomp_task_icv *icv = gomp_icv (true); - tgt->prev = icv->target_data; - icv->target_data = tgt; - } + struct target_mem_desc *tgt + = gomp_map_vars (devicep, mapnum, hostaddrs, sizes, kinds, false); + struct gomp_task_icv *icv = gomp_icv (true); + tgt->prev = icv->target_data; + icv->target_data = tgt; } void @@ -483,15 +531,156 @@ void GOMP_target_update (int device, const void *openmp_target, size_t mapnum, void **hostaddrs, size_t *sizes, unsigned char *kinds) { - device = resolve_device (device); - if (device == -1) + struct gomp_device_descr *devicep = resolve_device (device); + if (devicep == NULL) return; - if (device == 257) - gomp_update (mapnum, hostaddrs, sizes, kinds); + gomp_update (devicep, mapnum, hostaddrs, sizes, kinds); } void GOMP_teams (unsigned int num_teams, unsigned int thread_limit) { } + +#ifdef PLUGIN_SUPPORT + +/* This function checks if the given string FNAME matches + "libgomp-plugin-*.so.1". */ +static bool +gomp_check_plugin_file_name (const char *fname) +{ + const char *prefix = "libgomp-plugin-"; + const char *suffix = ".so.1"; + if (!fname) + return false; + if (strncmp (fname, prefix, strlen (prefix)) != 0) + return false; + if (strncmp (fname + strlen (fname) - strlen (suffix), suffix, + strlen (suffix)) != 0) + return false; + return true; +} + +/* This function tries to load plugin for DEVICE. Name of plugin is passed + in PLUGIN_NAME. + Plugin handle and handles of the found functions are stored in the + corresponding fields of DEVICE. + The function returns TRUE on success and FALSE otherwise. */ +static bool +gomp_load_plugin_for_device (struct gomp_device_descr *device, + const char *plugin_name) +{ + if (!device || !plugin_name) + return false; + + device->plugin_handle = dlopen (plugin_name, RTLD_LAZY); + if (!device->plugin_handle) + return false; + + /* Clear any existing error. */ + dlerror (); + + /* Check if all required functions are available in the plugin and store + their handlers. + TODO: check for other routines as well. */ + device->device_available_func = dlsym (device->plugin_handle, + "device_available"); + if (dlerror () != NULL) + { + dlclose (device->plugin_handle); + return false; + } + + return true; +} + +/* This functions scans folder, specified in environment variable + LIBGOMP_PLUGIN_PATH, and loads all suitable libgomp plugins from this folder. + For a plugin to be suitable, its name should be "libgomp-plugin-*.so.1" and + it should implement a certain set of functions. + Result of this function is properly initialized variable NUM_DEVICES and + array DEVICES, containing all plugins and their callback handles. */ +static void +gomp_find_available_plugins (void) +{ + char *plugin_path = NULL; + DIR *dir = NULL; + struct dirent *ent; + char plugin_name[PATH_MAX]; + + num_devices = 0; + devices = NULL; + + plugin_path = getenv ("LIBGOMP_PLUGIN_PATH"); + if (!plugin_path) + return; + + dir = opendir (plugin_path); + if (!dir) + return; + + while ((ent = readdir (dir)) != NULL) + { + struct gomp_device_descr current_device; + if (!gomp_check_plugin_file_name (ent->d_name)) + continue; + if (strlen (plugin_path) + 1 + strlen (ent->d_name) >= PATH_MAX) + continue; + strcpy (plugin_name, plugin_path); + strcat (plugin_name, "/"); + strcat (plugin_name, ent->d_name); + if (!gomp_load_plugin_for_device (¤t_device, plugin_name)) + continue; + devices = realloc (devices, (num_devices + 1) + * sizeof (struct gomp_device_descr)); + if (devices == NULL) + { + num_devices = 0; + closedir (dir); + return; + } + + devices[num_devices] = current_device; + devices[num_devices].id = num_devices + 1; + devices[num_devices].dev_splay_tree.root = NULL; + gomp_mutex_init (&devices[num_devices].dev_env_lock); + num_devices++; + } + closedir (dir); + + /* FIXME: Temporary hack for testing non-shared address spaces on host. + We create device 257 just to check memory mapping. */ + if (num_devices == 0) + { + num_devices = 1; + devices = malloc (sizeof (struct gomp_device_descr)); + if (devices == NULL) + { + num_devices = 0; + return; + } + devices[0].plugin_handle = NULL; + devices[0].device_available_func = NULL; + devices[0].dev_splay_tree.root = NULL; + gomp_mutex_init (&devices[0].dev_env_lock); + } + devices[0].id = 257; +} + +/* This function initializes runtime needed for offloading. + It loads plugins, sets up a connection with devices, etc. */ +static void +gomp_target_init (void) +{ + gomp_find_available_plugins (); +} + +#else /* PLUGIN_SUPPORT */ +/* If dlfcn.h is unavailable we always fallback to host execution. + GOMP_target* routines are just stubs for this case. */ +static void +gomp_target_init (void) +{ +} +#endif /* PLUGIN_SUPPORT */