@@ -30,6 +30,9 @@
/* Define to 1 if you have the <inttypes.h> 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 <memory.h> 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
@@ -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 :
@@ -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)
@@ -30,6 +30,15 @@
#include <stdlib.h>
#include <string.h>
+#ifdef PLUGIN_SUPPORT
+# include <dlfcn.h>
+# include <dirent.h>
+#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,56 @@ 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 name. */
+ char plugin_name[PATH_MAX];
+
+ /* 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;
+ return num_devices;
}
-static int
-resolve_device (int device)
+static struct gomp_device_descr*
+resolve_device (int device_id)
{
- if (device == -1)
+ (void) pthread_once (&gomp_is_initialized, gomp_target_init);
+ 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];
+
+ if (device_id >= gomp_get_num_devices ())
+ return NULL;
+ 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,15 +179,20 @@ 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;
- 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);
tgt->list_count = mapnum;
tgt->refcount = 1;
+ tgt->device_descr = devicep;
+
+ if (!devicep)
+ return tgt;
if (mapnum == 0)
return tgt;
@@ -159,7 +206,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 +214,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 +263,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 +283,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 +305,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 +354,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 +373,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 +382,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 +394,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 +405,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 +457,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
@@ -417,17 +474,20 @@ GOMP_target (int device, void (*fn) (void *), const char *fnname,
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)
+ /* 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. */
+ 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);
gomp_unmap_vars (tgt);
}
@@ -437,8 +497,8 @@ void
GOMP_target_data (int device, 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);
@@ -449,17 +509,17 @@ GOMP_target_data (int device, size_t mapnum, void **hostaddrs, size_t *sizes,
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 (devicep, 0, NULL, NULL, NULL, false);
tgt->prev = icv->target_data;
icv->target_data = tgt;
}
return;
}
- if (device == 257)
+ if (devicep->id == 257)
{
struct target_mem_desc *tgt
- = gomp_map_vars (mapnum, hostaddrs, sizes, kinds, false);
+ = 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;
@@ -482,15 +542,149 @@ void
GOMP_target_update (int device, 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);
+ if (devicep->id == 257)
+ 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 should be
+ stored in PLUGIN_NAME field.
+ 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)
+{
+ if (!device || !device->plugin_name)
+ return false;
+
+ device->plugin_handle = dlopen (device->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;
+
+ 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;
+ strncpy (current_device.plugin_name, plugin_path, PATH_MAX);
+ strcat (current_device.plugin_name, "/");
+ strcat (current_device.plugin_name, ent->d_name);
+ if (!gomp_load_plugin_for_device (¤t_device))
+ 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;
+ 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].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 */