This is the mail archive of the gcc-patches@gcc.gnu.org mailing list for the GCC project.


Index Nav: [Date Index] [Subject Index] [Author Index] [Thread Index]
Message Nav: [Date Prev] [Date Next] [Thread Prev] [Thread Next]
Other format: [Raw text]

[PATCH][gomp4] Plugins Support in LibGOMP (Take 2)


Hi Jakub,

I merged my patch with recent changes in gomp4-branch, and the new version is
below.  Also, I fixed most of your remarks - the one isn't fixed is checking
sizeof(void*)==sizeof(uintptr_t) in configure.  I'll do it in the next patch.

Is it ok for gomp4-branch?

Also, I was thinking of how we could test such changes.  Maybe we need to add
some logging stuff to all these libgomp routines - then we could check logs and
thus test if everything works as expected.  Otherwise it's really hard to find
out, whether offloading happened or not.  What do you think?

Thanks, Michael

---
 libgomp/config.h.in  |   6 ++
 libgomp/configure    |  63 +++++++++++
 libgomp/configure.ac |   9 ++
 libgomp/target.c     | 290 ++++++++++++++++++++++++++++++++++++++++++---------
 4 files changed, 320 insertions(+), 48 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 <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
 
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 e9b3386..fd2383e 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -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 (&current_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 */
-- 
1.8.3.1


Index Nav: [Date Index] [Subject Index] [Author Index] [Thread Index]
Message Nav: [Date Prev] [Date Next] [Thread Prev] [Thread Next]