diff mbox

Fwd: [RFC][gomp4] Offloading patches (2/3): Add tables generation

Message ID 20140308145015.GA61489@msticlxl57.ims.intel.com
State New
Headers show

Commit Message

Ilya Verbin March 8, 2014, 2:50 p.m. UTC 
Hi Bernd,

Here is updated patch for libgomp.  It assumes that there is a constructor with
a call to GOMP_offload_register in every target image, created by mkoffload
tool.  How does this look?

---
 libgomp/libgomp.map   |    1 +
 libgomp/plugin-host.c |   58 ++++++++++++++++-
 libgomp/target.c      |  170 +++++++++++++++++++++++++++++++++++++++++++++----
 3 files changed, 213 insertions(+), 16 deletions(-)

Comments

Bernd Schmidt March 12, 2014, 2:12 p.m. UTC | #1 
Hi,

On 03/08/2014 03:50 PM, Ilya Verbin wrote:
> Here is updated patch for libgomp.  It assumes that there is a constructor with
> a call to GOMP_offload_register in every target image, created by mkoffload
> tool.  How does this look?

LGTM. Shall I start committing my changes to the branch?


Bernd
Ilya Verbin March 12, 2014, 2:51 p.m. UTC | #2 
2014-03-12 18:12 GMT+04:00 Bernd Schmidt <bernds@codesourcery.com>:
> LGTM. Shall I start committing my changes to the branch?

Yes, I think you should commit your changes.
And we will rewrite our part to use the new configure approach.

  -- Ilya
Thomas Schwinge March 17, 2014, 3 p.m. UTC | #3 
Hi!

On Sat, 8 Mar 2014 18:50:15 +0400, Ilya Verbin <iverbin@gmail.com> wrote:
> --- a/libgomp/libgomp.map
> +++ b/libgomp/libgomp.map
> @@ -208,6 +208,7 @@ GOMP_3.0 {
>  
>  GOMP_4.0 {
>    global:
> +	GOMP_offload_register;
>  	GOMP_barrier_cancel;
>  	GOMP_cancel;
>  	GOMP_cancellation_point;

Now that the GOMP_4.0 symbol version is being used in GCC trunk, and will
be in the GCC 4.9 release, can we still add new symbols to it here?
(Jakub?)

> --- a/libgomp/plugin-host.c
> +++ b/libgomp/plugin-host.c

> +const int TARGET_TYPE_HOST = 0;

We'll have to see whether this (that is, libgomp/target.c:enum
target_type) should live in a shared header file, but OK for the moment.

> +void
> +device_run (void *fn_ptr, void *vars)
> +{
> +#ifdef DEBUG
> +  printf ("libgomp plugin: %s:%s (%p, %p)\n", __FILE__, __FUNCTION__, fn_ptr,
> +	  vars);
> +#endif
> +
> +  void (*fn)(void *) = (void (*)(void *)) fn_ptr;
> +
> +  fn (vars);
> +}

Why not make fn_ptr a proper function pointer?  Ah, because of
GOMP_target passing (void *) tgt_fn->tgt->tgt_start for the
!TARGET_TYPE_HOST case...

Would it make sense to have device_run return a value to make it able to
indicate to libgomp that the function cannot be run on the device (for
whatever reason), and libgomp should use host-fallback execution?
(Probably that needs more thought and discussion, OK to defer.)

> --- a/libgomp/target.c
> +++ b/libgomp/target.c

> +enum target_type {
> +  TARGET_TYPE_HOST,
> +  TARGET_TYPE_INTEL_MIC
> +};

(As discussed above, but OK to defer.)

> @@ -120,15 +140,26 @@ struct gomp_device_descr
>       TARGET construct.  */
>    int id;
>  
> +  /* This is the TYPE of device.  */
> +  int type;

Use enum target_type instead of int?

> +/* This function should be called from every offload image.  It gets the
> +   descriptor of the host func and var tables HOST_TABLE, TYPE of the target,
> +   and TARGET_DATA needed by target plugin (target tables, etc.)  */
> +void
> +GOMP_offload_register (void *host_table, int type, void *target_data)
> +{
> +  offload_images = realloc (offload_images,
> +			    (num_offload_images + 1)
> +			    * sizeof (struct offload_image_descr));
> +
> +  if (offload_images == NULL)
> +    return;

Fail silently, or use gomp_realloc to fail loudly?

> @@ -701,16 +836,25 @@ gomp_find_available_plugins (void)

> - out:
> +out:

Emacs wants the space to be there, so I assume that's the coding standard
to use.  ;-)

>    if (dir)
>      closedir (dir);
> +  free (offload_images);

I suggest to set offload_images = NULL, for clarity.

> +  num_offload_images = 0;
>  }

We may need to revisit this later: currently it's not possible to
register additional plugins after libgomp has initialized
(gomp_target_init, gomp_find_available_plugins just executed once), but
should that ever be made possible, we'd need to preserve offload_images.


OK to commit, thanks!


Grüße,
 Thomas
Jakub Jelinek March 17, 2014, 4:47 p.m. UTC | #4 
On Mon, Mar 17, 2014 at 04:00:11PM +0100, Thomas Schwinge wrote:
> Hi!
> 
> On Sat, 8 Mar 2014 18:50:15 +0400, Ilya Verbin <iverbin@gmail.com> wrote:
> > --- a/libgomp/libgomp.map
> > +++ b/libgomp/libgomp.map
> > @@ -208,6 +208,7 @@ GOMP_3.0 {
> >  
> >  GOMP_4.0 {
> >    global:
> > +	GOMP_offload_register;
> >  	GOMP_barrier_cancel;
> >  	GOMP_cancel;
> >  	GOMP_cancellation_point;
> 
> Now that the GOMP_4.0 symbol version is being used in GCC trunk, and will
> be in the GCC 4.9 release, can we still add new symbols to it here?
> (Jakub?)

If GCC 4.9 release will not include that symbol, then it must be in a new
symbol version, e.g. GOMP_4.1 (note, the fact that GOMP_ symbol version
matched now the OpenMP standard version wasn't always true and might not be
true always either (or we could use GOMP_4.0.1 symver).

	Jakub
Ilya Verbin March 18, 2014, 5:16 p.m. UTC | #5 
On 17 Mar 16:00, Thomas Schwinge wrote:
> >  GOMP_4.0 {
> >    global:
> > +	GOMP_offload_register;
> >  	GOMP_barrier_cancel;
> >  	GOMP_cancel;
> >  	GOMP_cancellation_point;
> 
> Now that the GOMP_4.0 symbol version is being used in GCC trunk, and will
> be in the GCC 4.9 release, can we still add new symbols to it here?
> (Jakub?)

I moved it to GOMP_4.0.1.

> > +  /* This is the TYPE of device.  */
> > +  int type;
> 
> Use enum target_type instead of int?

Done.

> > +  offload_images = realloc (offload_images,
> > +			    (num_offload_images + 1)
> > +			    * sizeof (struct offload_image_descr));
> > +
> > +  if (offload_images == NULL)
> > +    return;
> 
> Fail silently, or use gomp_realloc to fail loudly?

Replaced with gomp_realloc.

> >    if (dir)
> >      closedir (dir);
> > +  free (offload_images);
> 
> I suggest to set offload_images = NULL, for clarity.

Done.

> OK to commit, thanks!

Committed as r208657.


> Would it make sense to have device_run return a value to make it able to
> indicate to libgomp that the function cannot be run on the device (for
> whatever reason), and libgomp should use host-fallback execution?
> (Probably that needs more thought and discussion, OK to defer.)

Consider the following example (using OpenMP, I don't know OpenACC :)

int foo ()
{
  int x = 0;

  /* offload_fn1  */
  #pragma omp target map(to: x)
    {
      x += 5;
    }

  /* Some code on host without updating 'x' from target.  */

  /* offload_fn2  */
  #pragma omp target map(from: x)
    {
      x += 10;
    }

  return x;
}

If both offload_fn1 and offload_fn2 are executed on host, everything is fine
and x = 15.  The same goes to the case when both offload_fn1 and offload_fn2
are executed on target.  But if offload_fn1 is executed on target and
offload_fn2 is executed on host, then 'x' will have incorrect value (10).

Therefore, I proposed to check for target device availability only during
initialization of the plugin.  And to make a decision at this point, will
libgomp run all functions on host or on target.  Probably libgomp should return
an error if something was executed on device, but then it becomes unavailable.

  -- Ilya
Bernd Schmidt March 20, 2014, 5:09 p.m. UTC | #6 
On 03/12/2014 03:51 PM, Ilya Verbin wrote:
> 2014-03-12 18:12 GMT+04:00 Bernd Schmidt <bernds@codesourcery.com>:
>> LGTM. Shall I start committing my changes to the branch?
>
> Yes, I think you should commit your changes.
> And we will rewrite our part to use the new configure approach.

Done now.  I think/hope that I've committed all the ones that are not 
entirely related to ptx - let me know if you find there are any missing 
pieces.


Bernd
diff mbox

Patch

diff --git a/libgomp/libgomp.map b/libgomp/libgomp.map
index d8631a6..e43cb42 100644
--- a/libgomp/libgomp.map
+++ b/libgomp/libgomp.map
@@ -208,6 +208,7 @@  GOMP_3.0 {
 
 GOMP_4.0 {
   global:
+	GOMP_offload_register;
 	GOMP_barrier_cancel;
 	GOMP_cancel;
 	GOMP_cancellation_point;
diff --git a/libgomp/plugin-host.c b/libgomp/plugin-host.c
index 5354ebe..ec0c78c 100644
--- a/libgomp/plugin-host.c
+++ b/libgomp/plugin-host.c
@@ -33,14 +33,53 @@ 
 #include <stdlib.h>
 #include <string.h>
 
-bool
-device_available (void)
+const int TARGET_TYPE_HOST = 0;
+
+int
+get_type (void)
 {
 #ifdef DEBUG
   printf ("libgomp plugin: %s:%s\n", __FILE__, __FUNCTION__);
 #endif
 
-  return true;
+  return TARGET_TYPE_HOST;
+}
+
+int
+get_num_devices (void)
+{
+#ifdef DEBUG
+  printf ("libgomp plugin: %s:%s\n", __FILE__, __FUNCTION__);
+#endif
+
+  return 1;
+}
+
+void
+offload_register (void *host_table, void *target_data)
+{
+#ifdef DEBUG
+  printf ("libgomp plugin: %s:%s (%p, %p)\n", __FILE__, __FUNCTION__,
+	  host_table, target_data);
+#endif
+}
+
+void
+device_init (void)
+{
+#ifdef DEBUG
+  printf ("libgomp plugin: %s:%s\n", __FILE__, __FUNCTION__);
+#endif
+}
+
+int
+device_get_table (void *table)
+{
+#ifdef DEBUG
+  printf ("libgomp plugin: %s:%s (%p)\n", __FILE__, __FUNCTION__, table);
+#endif
+
+  return 0;
 }
 
 void *
@@ -82,3 +121,16 @@  void *device_host2dev (void *dest, const void *src, size_t n)
 
   return memcpy (dest, src, n);
 }
+
+void
+device_run (void *fn_ptr, void *vars)
+{
+#ifdef DEBUG
+  printf ("libgomp plugin: %s:%s (%p, %p)\n", __FILE__, __FUNCTION__, fn_ptr,
+	  vars);
+#endif
+
+  void (*fn)(void *) = (void (*)(void *)) fn_ptr;
+
+  fn (vars);
+}
diff --git a/libgomp/target.c b/libgomp/target.c
index dbe6e28..8be9ea1 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -87,6 +87,26 @@  struct splay_tree_key_s {
   bool copy_from;
 };
 
+enum target_type {
+  TARGET_TYPE_HOST,
+  TARGET_TYPE_INTEL_MIC
+};
+
+/* This structure describes an offload image.
+   It contains type of the target, pointer to host table descriptor, and pointer
+   to target data.  */
+struct offload_image_descr {
+  int type;
+  void *host_table;
+  void *target_data;
+};
+
+/* Array of descriptors of offload images.  */
+static struct offload_image_descr *offload_images;
+
+/* Total number of offload images.  */
+static int num_offload_images;
+
 /* Array of descriptors of all available devices.  */
 static struct gomp_device_descr *devices;
 
@@ -120,15 +140,26 @@  struct gomp_device_descr
      TARGET construct.  */
   int id;
 
+  /* This is the TYPE of device.  */
+  int type;
+
+  /* Set to true when device is initialized.  */
+  bool is_initialized;
+
   /* Plugin file handler.  */
   void *plugin_handle;
 
   /* Function handlers.  */
-  bool (*device_available_func) (void);
+  int (*get_type_func) (void);
+  int (*get_num_devices_func) (void);
+  void (*offload_register_func) (void *, void *);
+  void (*device_init_func) (void);
+  int (*device_get_table_func) (void *);
   void *(*device_alloc_func) (size_t);
   void (*device_free_func) (void *);
-  void *(*device_dev2host_func)(void *, const void *, size_t);
-  void *(*device_host2dev_func)(void *, const void *, size_t);
+  void *(*device_dev2host_func) (void *, const void *, size_t);
+  void *(*device_host2dev_func) (void *, const void *, size_t);
+  void (*device_run_func) (void *, void *);
 
   /* Splay tree containing information about mapped memory regions.  */
   struct splay_tree_s dev_splay_tree;
@@ -137,6 +168,13 @@  struct gomp_device_descr
   gomp_mutex_t dev_env_lock;
 };
 
+struct mapping_table {
+  uintptr_t host_start;
+  uintptr_t host_end;
+  uintptr_t tgt_start;
+  uintptr_t tgt_end;
+};
+
 attribute_hidden int
 gomp_get_num_devices (void)
 {
@@ -474,6 +512,63 @@  gomp_update (struct gomp_device_descr *devicep, size_t mapnum,
   gomp_mutex_unlock (&devicep->dev_env_lock);
 }
 
+/* This function should be called from every offload image.  It gets the
+   descriptor of the host func and var tables HOST_TABLE, TYPE of the target,
+   and TARGET_DATA needed by target plugin (target tables, etc.)  */
+void
+GOMP_offload_register (void *host_table, int type, void *target_data)
+{
+  offload_images = realloc (offload_images,
+			    (num_offload_images + 1)
+			    * sizeof (struct offload_image_descr));
+
+  if (offload_images == NULL)
+    return;
+
+  offload_images[num_offload_images].type = type;
+  offload_images[num_offload_images].host_table = host_table;
+  offload_images[num_offload_images].target_data = target_data;
+
+  num_offload_images++;
+}
+
+static void
+gomp_init_device (struct gomp_device_descr *devicep)
+{
+  /* Initialize the target device.  */
+  devicep->device_init_func ();
+
+  /* Get address mapping table for device.  */
+  struct mapping_table *table = NULL;
+  int num_entries = devicep->device_get_table_func (&table);
+
+  /* Insert host-target address mapping into dev_splay_tree.  */
+  int i;
+  for (i = 0; i < num_entries; i++)
+    {
+      struct target_mem_desc *tgt = gomp_malloc (sizeof (*tgt));
+      tgt->refcount = 1;
+      tgt->array = gomp_malloc (sizeof (*tgt->array));
+      tgt->tgt_start = table[i].tgt_start;
+      tgt->tgt_end = table[i].tgt_end;
+      tgt->to_free = NULL;
+      tgt->list_count = 0;
+      tgt->device_descr = devicep;
+      splay_tree_node node = tgt->array;
+      splay_tree_key k = &node->key;
+      k->host_start = table[i].host_start;
+      k->host_end = table[i].host_end;
+      k->tgt_offset = 0;
+      k->tgt = tgt;
+      node->left = NULL;
+      node->right = NULL;
+      splay_tree_insert (&devicep->dev_splay_tree, node);
+    }
+
+  free (table);
+  devicep->is_initialized = true;
+}
+
 /* Called when encountering a target directive.  If DEVICE
    is -1, it means use device-var ICV.  If it is -2 (or any other value
    larger than last available hw device, use host fallback.
@@ -507,7 +602,17 @@  GOMP_target (int device, void (*fn) (void *), const void *openmp_target,
       return;
     }
 
-  struct target_mem_desc *tgt
+  if (!devicep->is_initialized)
+    gomp_init_device (devicep);
+
+  struct splay_tree_key_s k;
+  k.host_start = (uintptr_t) fn;
+  k.host_end = k.host_start + 1;
+  splay_tree_key tgt_fn = splay_tree_lookup (&devicep->dev_splay_tree, &k);
+  if (tgt_fn == NULL && devicep->type != TARGET_TYPE_HOST)
+    gomp_fatal ("Target function wasn't mapped");
+
+  struct target_mem_desc *tgt_vars
     = gomp_map_vars (devicep, mapnum, hostaddrs, sizes, kinds, true);
   struct gomp_thread old_thr, *thr = gomp_thread ();
   old_thr = *thr;
@@ -517,10 +622,14 @@  GOMP_target (int device, void (*fn) (void *), const void *openmp_target,
       thr->place = old_thr.place;
       thr->ts.place_partition_len = gomp_places_list_len;
     }
-  fn ((void *) tgt->tgt_start);
+  if (devicep->type == TARGET_TYPE_HOST)
+    devicep->device_run_func (fn, (void *) tgt_vars->tgt_start);
+  else
+    devicep->device_run_func ((void *) tgt_fn->tgt->tgt_start,
+			      (void *) tgt_vars->tgt_start);
   gomp_free_thread (thr);
   *thr = old_thr;
-  gomp_unmap_vars (tgt);
+  gomp_unmap_vars (tgt_vars);
 }
 
 void
@@ -546,6 +655,9 @@  GOMP_target_data (int device, const void *openmp_target, size_t mapnum,
       return;
     }
 
+  if (!devicep->is_initialized)
+    gomp_init_device (devicep);
+
   struct target_mem_desc *tgt
     = gomp_map_vars (devicep, mapnum, hostaddrs, sizes, kinds, false);
   struct gomp_task_icv *icv = gomp_icv (true);
@@ -573,6 +685,9 @@  GOMP_target_update (int device, const void *openmp_target, size_t mapnum,
   if (devicep == NULL)
     return;
 
+  if (!devicep->is_initialized)
+    gomp_init_device (devicep);
+
   gomp_update (devicep, mapnum, hostaddrs, sizes, kinds);
 }
 
@@ -639,11 +754,16 @@  gomp_load_plugin_for_device (struct gomp_device_descr *device,
 	goto out;							\
     }									\
   while (0)
-  DLSYM (device_available);
+  DLSYM (get_type);
+  DLSYM (get_num_devices);
+  DLSYM (offload_register);
+  DLSYM (device_init);
+  DLSYM (device_get_table);
   DLSYM (device_alloc);
   DLSYM (device_free);
   DLSYM (device_dev2host);
   DLSYM (device_host2dev);
+  DLSYM (device_run);
 #undef DLSYM
 
  out:
@@ -656,6 +776,21 @@  gomp_load_plugin_for_device (struct gomp_device_descr *device,
   return err == NULL;
 }
 
+/* This function finds OFFLOAD_IMAGES corresponding to DEVICE type, and
+   registers them in the plugin.  */
+static void
+gomp_register_images_for_device (struct gomp_device_descr *device)
+{
+  int i;
+  for (i = 0; i < num_offload_images; i++)
+    {
+      struct offload_image_descr *image = &offload_images[i];
+
+      if (device->type == image->type || device->type == TARGET_TYPE_HOST)
+	device->offload_register_func (image->host_table, image->target_data);
+    }
+}
+
 /* 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
@@ -701,16 +836,25 @@  gomp_find_available_plugins (void)
 	  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++;
+      /* FIXME: Properly handle multiple devices of the same type.  */
+      if (current_device.get_num_devices_func () >= 1)
+	{
+	  current_device.id = num_devices + 1;
+	  current_device.type = current_device.get_type_func ();
+	  current_device.is_initialized = false;
+	  current_device.dev_splay_tree.root = NULL;
+	  gomp_register_images_for_device (&current_device);
+	  devices[num_devices] = current_device;
+	  gomp_mutex_init (&devices[num_devices].dev_env_lock);
+	  num_devices++;
+	}
     }
 
- out:
+out:
   if (dir)
     closedir (dir);
+  free (offload_images);
+  num_offload_images = 0;
 }
 
 /* This function initializes runtime needed for offloading.