From patchwork Tue Sep 23 18:17:50 2014 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Julian Brown X-Patchwork-Id: 392614 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.2 with cipher ECDHE-RSA-AES256-GCM-SHA384 (256/256 bits)) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id EC240140082 for ; Wed, 24 Sep 2014 04:18:16 +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:subject:message-id:mime-version:content-type; q=dns; s= default; b=X7ps6atr5qh87xwUtfJmkSJ8YpiEMqKszgjayPzP09oEEQl5irzXC 9esdjPYlPk+lqjcy8eX6PGVZJvzbvJZz/hTXE9xnu7Z/VS42W/5l1I6UuIBDL0eh YvKozQ6iIO4USG+eCp2+U+bFdVznKs7RXuoXAlxOVZuZ8VhEx20Jn4= 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:subject:message-id:mime-version:content-type; s= default; bh=gxmnp2mvXjuCjDd3ZNI4FzPYeNk=; b=vqUG7tYJyji93k8QhUkR L/vfOfp3YE4mIP0lORfOpc3a0JJFtc8knE/yfSXTIASR0emBTxDZVPtSoEEUAUYk r+hPBW1Sgn25u7C7waFHE9jEXf++Vo6p4vELFFz0S0CHfzj/ivHVCRzduttvA9U6 yBKXv0/+y2xf1s2OpjYxSHs= Received: (qmail 13478 invoked by alias); 23 Sep 2014 18:18:07 -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 13354 invoked by uid 89); 23 Sep 2014 18:18:06 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-1.9 required=5.0 tests=AWL, BAYES_00 autolearn=ham version=3.3.2 X-HELO: relay1.mentorg.com Received: from relay1.mentorg.com (HELO relay1.mentorg.com) (192.94.38.131) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Tue, 23 Sep 2014 18:18:01 +0000 Received: from nat-ies.mentorg.com ([192.94.31.2] helo=SVR-IES-FEM-01.mgc.mentorg.com) by relay1.mentorg.com with esmtp id 1XWUez-0002bt-E5 from Julian_Brown@mentor.com for gcc-patches@gcc.gnu.org; Tue, 23 Sep 2014 11:17:58 -0700 Received: from octopus (137.202.0.76) by SVR-IES-FEM-01.mgc.mentorg.com (137.202.0.104) with Microsoft SMTP Server id 14.3.181.6; Tue, 23 Sep 2014 19:17:55 +0100 Date: Tue, 23 Sep 2014 19:17:50 +0100 From: Julian Brown To: Subject: [PATCH 2/10] OpenACC 2.0 support for libgomp - initial plugin support Message-ID: <20140923191750.5c1c56ce@octopus> MIME-Version: 1.0 X-IsSubscribed: yes This patch is by Michael Zolotukhin and was originally posted here: https://gcc.gnu.org/ml/gcc-patches/2013-09/msg01469.html It contains an initial implementation of plugin support for libgomp, for implementing different hardware devices for pieces of accelerated code to be offloaded to. I also merged a minor follow-up fix by Thomas Schwinge. Julian xxxx-xx-xx Michael Zolotukhin Thomas Schwinge * configure.ac: Add checks for plugins support. * config.h.in: Regenerated. * configure: Regenerated. * target.c (struct target_mem_desc): Add device_descr field. (devices): New. (num_devices): New. (struct gomp_device_descr): New. (gomp_get_num_devices): Call gomp_target_init. (resolve_device): Return device_descr instead of int. (gomp_map_vars): Add devicep argument and update the function accordingly. (gomp_unmap_tgt): Likewise. (gomp_unmap_vars): Likewise. (gomp_update): Likewise. (GOMP_target): Use device_descr struct. (GOMP_target_data): Likewise. (GOMP_target_update): Likewise. (gomp_check_plugin_file_name): New. (gomp_load_plugin_for_device): New. (gomp_find_available_plugins): New. (gomp_target_init): New. commit 75ef137a74cbd6af36a75b30edf60350ec9eae0d Author: Julian Brown Date: Fri Sep 19 08:51:44 2014 -0700 Merge of r202827. xxxx-xx-xx Michael Zolotukhin Thomas Schwinge * configure.ac: Add checks for plugins support. * config.h.in: Regenerated. * configure: Regenerated. * target.c (struct target_mem_desc): Add device_descr field. (devices): New. (num_devices): New. (struct gomp_device_descr): New. (gomp_get_num_devices): Call gomp_target_init. (resolve_device): Return device_descr instead of int. (gomp_map_vars): Add devicep argument and update the function accordingly. (gomp_unmap_tgt): Likewise. (gomp_unmap_vars): Likewise. (gomp_update): Likewise. (GOMP_target): Use device_descr struct. (GOMP_target_data): Likewise. (GOMP_target_update): Likewise. (gomp_check_plugin_file_name): New. (gomp_load_plugin_for_device): New. (gomp_find_available_plugins): New. (gomp_target_init): New. 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 766eb09..704f22a 100755 --- a/libgomp/configure +++ b/libgomp/configure @@ -15052,6 +15052,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 84d250f..da06426 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 9b759b1..55b3781 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -31,6 +31,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; @@ -51,6 +60,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[]; @@ -71,6 +84,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 @@ -88,33 +107,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. */ @@ -138,8 +179,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; @@ -147,6 +189,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_align = align; 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); } @@ -419,8 +472,8 @@ 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. */ struct gomp_thread old_thr, *thr = gomp_thread (); @@ -436,21 +489,19 @@ GOMP_target (int device, void (*fn) (void *), const void *openmp_target, *thr = old_thr; 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); @@ -461,21 +512,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 @@ -494,12 +542,11 @@ 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 @@ -513,3 +560,146 @@ GOMP_teams (unsigned int num_teams, unsigned int thread_limit) } (void) num_teams; } + +#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) + goto out; + + dir = opendir (plugin_path); + if (!dir) + goto out; + + 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); + goto out; + } + + 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); + + out: + /* 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 */