libgomp/target.c

   1 /* Copyright (C) 2013-2014 Free Software Foundation, Inc.
   2    Contributed by Jakub Jelinek <jakub@redhat.com>.
   3
   4    This file is part of the GNU OpenMP Library (libgomp).
   5
   6    Libgomp is free software; you can redistribute it and/or modify it
   7    under the terms of the GNU General Public License as published by
   8    the Free Software Foundation; either version 3, or (at your option)
   9    any later version.
  10
  11    Libgomp is distributed in the hope that it will be useful, but WITHOUT ANY
  12    WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS
  13    FOR A PARTICULAR PURPOSE.  See the GNU General Public License for
  14    more details.
  15
  16    Under Section 7 of GPL version 3, you are granted additional
  17    permissions described in the GCC Runtime Library Exception, version
  18    3.1, as published by the Free Software Foundation.
  19
  20    You should have received a copy of the GNU General Public License and
  21    a copy of the GCC Runtime Library Exception along with this program;
  22    see the files COPYING3 and COPYING.RUNTIME respectively.  If not, see
  23    <http://www.gnu.org/licenses/>.  */
  24
  25 /* This file contains the support of offloading.  */
  26
  27 #include "config.h"
  28 #include "libgomp.h"
  29 #include "libgomp_target.h"
  30 #include <limits.h>
  31 #include <stdbool.h>
  32 #include <stdlib.h>
  33 #include <string.h>
  34
  35 #ifdef PLUGIN_SUPPORT
  36 #include <dlfcn.h>
  37 #endif
  38
  39 static void gomp_target_init (void);
  40
  41 static pthread_once_t gomp_is_initialized = PTHREAD_ONCE_INIT;
  42
  43 /* Forward declaration for a node in the tree.  */
  44 typedef struct splay_tree_node_s *splay_tree_node;
  45 typedef struct splay_tree_s *splay_tree;
  46 typedef struct splay_tree_key_s *splay_tree_key;
  47
  48 struct target_mem_desc {
  49   /* Reference count.  */
  50   uintptr_t refcount;
  51   /* All the splay nodes allocated together.  */
  52   splay_tree_node array;
  53   /* Start of the target region.  */
  54   uintptr_t tgt_start;
  55   /* End of the targer region.  */
  56   uintptr_t tgt_end;
  57   /* Handle to free.  */
  58   void *to_free;
  59   /* Previous target_mem_desc.  */
  60   struct target_mem_desc *prev;
  61   /* Number of items in following list.  */
  62   size_t list_count;
  63
  64   /* Corresponding target device descriptor.  */
  65   struct gomp_device_descr *device_descr;
  66
  67   /* List of splay keys to remove (or decrease refcount)
  68      at the end of region.  */
  69   splay_tree_key list[];
  70 };
  71
  72 struct splay_tree_key_s {
  73   /* Address of the host object.  */
  74   uintptr_t host_start;
  75   /* Address immediately after the host object.  */
  76   uintptr_t host_end;
  77   /* Descriptor of the target memory.  */
  78   struct target_mem_desc *tgt;
  79   /* Offset from tgt->tgt_start to the start of the target object.  */
  80   uintptr_t tgt_offset;
  81   /* Reference count.  */
  82   uintptr_t refcount;
  83   /* True if data should be copied from device to host at the end.  */
  84   bool copy_from;
  85 };
  86
  87 /* This structure describes an offload image.
  88    It contains type of the target device, pointer to host table descriptor, and
  89    pointer to target data.  */
  90 struct offload_image_descr {
  91   enum offload_target_type type;
  92   void *host_table;
  93   void *target_data;
  94 };
  95
  96 /* Array of descriptors of offload images.  */
  97 static struct offload_image_descr *offload_images;
  98
  99 /* Total number of offload images.  */
 100 static int num_offload_images;
 101
 102 /* Array of descriptors for all available devices.  */
 103 static struct gomp_device_descr *devices;
 104
 105 /* Total number of available devices.  */
 106 static int num_devices;
 107
 108 /* The comparison function.  */
 109
 110 static int
 111 splay_compare (splay_tree_key x, splay_tree_key y)
 112 {
 113   if (x->host_start == x->host_end
 114       && y->host_start == y->host_end)
 115     return 0;
 116   if (x->host_end <= y->host_start)
 117     return -1;
 118   if (x->host_start >= y->host_end)
 119     return 1;
 120   return 0;
 121 }
 122
 123 #include "splay-tree.h"
 124
 125 /* This structure describes accelerator device.
 126    It contains ID-number of the device, its type, function handlers for
 127    interaction with the device, and information about mapped memory.  */
 128 struct gomp_device_descr
 129 {
 130   /* This is the ID number of device.  It could be specified in DEVICE-clause of
 131      TARGET construct.  */
 132   int id;
 133
 134   /* This is the ID number of device among devices of the same type.  */
 135   int target_id;
 136
 137   /* This is the TYPE of device.  */
 138   enum offload_target_type type;
 139
 140   /* Set to true when device is initialized.  */
 141   bool is_initialized;
 142
 143   /* Function handlers.  */
 144   int (*get_type_func) (void);
 145   int (*get_num_devices_func) (void);
 146   void (*register_image_func) (void *, void *);
 147   void (*init_device_func) (int);
 148   int (*get_table_func) (int, void *);
 149   void *(*alloc_func) (int, size_t);
 150   void (*free_func) (int, void *);
 151   void *(*host2dev_func) (int, void *, const void *, size_t);
 152   void *(*dev2host_func) (int, void *, const void *, size_t);
 153   void (*run_func) (int, void *, void *);
 154
 155   /* Splay tree containing information about mapped memory regions.  */
 156   struct splay_tree_s dev_splay_tree;
 157
 158   /* Mutex for operating with the splay tree and other shared structures.  */
 159   gomp_mutex_t dev_env_lock;
 160 };
 161
 162 attribute_hidden int
 163 gomp_get_num_devices (void)
 164 {
 165   (void) pthread_once (&gomp_is_initialized, gomp_target_init);
 166   return num_devices;
 167 }
 168
 169 static struct gomp_device_descr *
 170 resolve_device (int device_id)
 171 {
 172   if (device_id == -1)
 173     {
 174       struct gomp_task_icv *icv = gomp_icv (false);
 175       device_id = icv->default_device_var;
 176     }
 177
 178   if (device_id < 0 || device_id >= gomp_get_num_devices ())
 179     return NULL;
 180
 181   return &devices[device_id];
 182 }
 183
 184
 185 /* Handle the case where splay_tree_lookup found oldn for newn.
 186    Helper function of gomp_map_vars.  */
 187
 188 static inline void
 189 gomp_map_vars_existing (splay_tree_key oldn, splay_tree_key newn,
 190                         unsigned char kind)
 191 {
 192   if (oldn->host_start > newn->host_start
 193       || oldn->host_end < newn->host_end)
 194     gomp_fatal ("Trying to map into device [%p..%p) object when"
 195                 "[%p..%p) is already mapped",
 196                 (void *) newn->host_start, (void *) newn->host_end,
 197                 (void *) oldn->host_start, (void *) oldn->host_end);
 198   oldn->refcount++;
 199 }
 200
 201 static struct target_mem_desc *
 202 gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
 203                void **hostaddrs, size_t *sizes, unsigned char *kinds,
 204                bool is_target)
 205 {
 206   size_t i, tgt_align, tgt_size, not_found_cnt = 0;
 207   struct splay_tree_key_s cur_node;
 208   struct target_mem_desc *tgt
 209     = gomp_malloc (sizeof (*tgt) + sizeof (tgt->list[0]) * mapnum);
 210   tgt->list_count = mapnum;
 211   tgt->refcount = 1;
 212   tgt->device_descr = devicep;
 213
 214   if (mapnum == 0)
 215     return tgt;
 216
 217   tgt_align = sizeof (void *);
 218   tgt_size = 0;
 219   if (is_target)
 220     {
 221       size_t align = 4 * sizeof (void *);
 222       tgt_align = align;
 223       tgt_size = mapnum * sizeof (void *);
 224     }
 225
 226   gomp_mutex_lock (&devicep->dev_env_lock);
 227   for (i = 0; i < mapnum; i++)
 228     {
 229       if (hostaddrs[i] == NULL)
 230         {
 231           tgt->list[i] = NULL;
 232           continue;
 233         }
 234       cur_node.host_start = (uintptr_t) hostaddrs[i];
 235       if ((kinds[i] & 7) != 4)
 236         cur_node.host_end = cur_node.host_start + sizes[i];
 237       else
 238         cur_node.host_end = cur_node.host_start + sizeof (void *);
 239       splay_tree_key n = splay_tree_lookup (&devicep->dev_splay_tree,
 240                                             &cur_node);
 241       if (n)
 242         {
 243           tgt->list[i] = n;
 244           gomp_map_vars_existing (n, &cur_node, kinds[i]);
 245         }
 246       else
 247         {
 248           size_t align = (size_t) 1 << (kinds[i] >> 3);
 249           tgt->list[i] = NULL;
 250           not_found_cnt++;
 251           if (tgt_align < align)
 252             tgt_align = align;
 253           tgt_size = (tgt_size + align - 1) & ~(align - 1);
 254           tgt_size += cur_node.host_end - cur_node.host_start;
 255           if ((kinds[i] & 7) == 5)
 256             {
 257               size_t j;
 258               for (j = i + 1; j < mapnum; j++)
 259                 if ((kinds[j] & 7) != 4)
 260                   break;
 261                 else if ((uintptr_t) hostaddrs[j] < cur_node.host_start
 262                          || ((uintptr_t) hostaddrs[j] + sizeof (void *)
 263                              > cur_node.host_end))
 264                   break;
 265                 else
 266                   {
 267                     tgt->list[j] = NULL;
 268                     i++;
 269                   }
 270             }
 271         }
 272     }
 273
 274   if (not_found_cnt || is_target)
 275     {
 276       /* Allocate tgt_align aligned tgt_size block of memory.  */
 277       /* FIXME: Perhaps change interface to allocate properly aligned
 278          memory.  */
 279       tgt->to_free = devicep->alloc_func (devicep->target_id,
 280                                           tgt_size + tgt_align - 1);
 281       tgt->tgt_start = (uintptr_t) tgt->to_free;
 282       tgt->tgt_start = (tgt->tgt_start + tgt_align - 1) & ~(tgt_align - 1);
 283       tgt->tgt_end = tgt->tgt_start + tgt_size;
 284     }
 285   else
 286     {
 287       tgt->to_free = NULL;
 288       tgt->tgt_start = 0;
 289       tgt->tgt_end = 0;
 290     }
 291
 292   tgt_size = 0;
 293   if (is_target)
 294     tgt_size = mapnum * sizeof (void *);
 295
 296   tgt->array = NULL;
 297   if (not_found_cnt)
 298     {
 299       tgt->array = gomp_malloc (not_found_cnt * sizeof (*tgt->array));
 300       splay_tree_node array = tgt->array;
 301       size_t j;
 302
 303       for (i = 0; i < mapnum; i++)
 304         if (tgt->list[i] == NULL)
 305           {
 306             if (hostaddrs[i] == NULL)
 307               continue;
 308             splay_tree_key k = &array->key;
 309             k->host_start = (uintptr_t) hostaddrs[i];
 310             if ((kinds[i] & 7) != 4)
 311               k->host_end = k->host_start + sizes[i];
 312             else
 313               k->host_end = k->host_start + sizeof (void *);
 314             splay_tree_key n
 315               = splay_tree_lookup (&devicep->dev_splay_tree, k);
 316             if (n)
 317               {
 318                 tgt->list[i] = n;
 319                 gomp_map_vars_existing (n, k, kinds[i]);
 320               }
 321             else
 322               {
 323                 size_t align = (size_t) 1 << (kinds[i] >> 3);
 324                 tgt->list[i] = k;
 325                 tgt_size = (tgt_size + align - 1) & ~(align - 1);
 326                 k->tgt = tgt;
 327                 k->tgt_offset = tgt_size;
 328                 tgt_size += k->host_end - k->host_start;
 329                 k->copy_from = false;
 330                 if ((kinds[i] & 7) == 2 || (kinds[i] & 7) == 3)
 331                   k->copy_from = true;
 332                 k->refcount = 1;
 333                 tgt->refcount++;
 334                 array->left = NULL;
 335                 array->right = NULL;
 336                 splay_tree_insert (&devicep->dev_splay_tree, array);
 337                 switch (kinds[i] & 7)
 338                   {
 339                   case 0: /* ALLOC */
 340                   case 2: /* FROM */
 341                     break;
 342                   case 1: /* TO */
 343                   case 3: /* TOFROM */
 344                     /* FIXME: Perhaps add some smarts, like if copying
 345                        several adjacent fields from host to target, use some
 346                        host buffer to avoid sending each var individually.  */
 347                     devicep->host2dev_func (devicep->target_id,
 348                                             (void *) (tgt->tgt_start
 349                                                       + k->tgt_offset),
 350                                             (void *) k->host_start,
 351                                             k->host_end - k->host_start);
 352                     break;
 353                   case 4: /* POINTER */
 354                     cur_node.host_start
 355                       = (uintptr_t) *(void **) k->host_start;
 356                     if (cur_node.host_start == (uintptr_t) NULL)
 357                       {
 358                         cur_node.tgt_offset = (uintptr_t) NULL;
 359                         devicep->host2dev_func (devicep->target_id,
 360                                                 (void *) (tgt->tgt_start
 361                                                           + k->tgt_offset),
 362                                                 (void *) &cur_node.tgt_offset,
 363                                                 sizeof (void *));
 364                         break;
 365                       }
 366                     /* Add bias to the pointer value.  */
 367                     cur_node.host_start += sizes[i];
 368                     cur_node.host_end = cur_node.host_start + 1;
 369                     n = splay_tree_lookup (&devicep->dev_splay_tree,
 370                                            &cur_node);
 371                     if (n == NULL)
 372                       {
 373                         /* Could be possibly zero size array section.  */
 374                         cur_node.host_end--;
 375                         n = splay_tree_lookup (&devicep->dev_splay_tree,
 376                                                &cur_node);
 377                         if (n == NULL)
 378                           {
 379                             cur_node.host_start--;
 380                             n = splay_tree_lookup (&devicep->dev_splay_tree,
 381                                                    &cur_node);
 382                             cur_node.host_start++;
 383                           }
 384                       }
 385                     if (n == NULL)
 386                       gomp_fatal ("Pointer target of array section "
 387                                   "wasn't mapped");
 388                     cur_node.host_start -= n->host_start;
 389                     cur_node.tgt_offset = n->tgt->tgt_start + n->tgt_offset
 390                                           + cur_node.host_start;
 391                     /* At this point tgt_offset is target address of the
 392                        array section.  Now subtract bias to get what we want
 393                        to initialize the pointer with.  */
 394                     cur_node.tgt_offset -= sizes[i];
 395                     devicep->host2dev_func (devicep->target_id,
 396                                             (void *) (tgt->tgt_start
 397                                                       + k->tgt_offset),
 398                                             (void *) &cur_node.tgt_offset,
 399                                             sizeof (void *));
 400                     break;
 401                   case 5: /* TO_PSET */
 402                     devicep->host2dev_func (devicep->target_id,
 403                                             (void *) (tgt->tgt_start
 404                                                       + k->tgt_offset),
 405                                             (void *) k->host_start,
 406                                             k->host_end - k->host_start);
 407                     for (j = i + 1; j < mapnum; j++)
 408                       if ((kinds[j] & 7) != 4)
 409                         break;
 410                       else if ((uintptr_t) hostaddrs[j] < k->host_start
 411                                || ((uintptr_t) hostaddrs[j] + sizeof (void *)
 412                                    > k->host_end))
 413                         break;
 414                       else
 415                         {
 416                           tgt->list[j] = k;
 417                           k->refcount++;
 418                           cur_node.host_start
 419                             = (uintptr_t) *(void **) hostaddrs[j];
 420                           if (cur_node.host_start == (uintptr_t) NULL)
 421                             {
 422                               cur_node.tgt_offset = (uintptr_t) NULL;
 423                               devicep->host2dev_func (devicep->target_id,
 424                                  (void *) (tgt->tgt_start + k->tgt_offset
 425                                            + ((uintptr_t) hostaddrs[j]
 426                                               - k->host_start)),
 427                                  (void *) &cur_node.tgt_offset,
 428                                  sizeof (void *));
 429                               i++;
 430                               continue;
 431                             }
 432                           /* Add bias to the pointer value.  */
 433                           cur_node.host_start += sizes[j];
 434                           cur_node.host_end = cur_node.host_start + 1;
 435                           n = splay_tree_lookup (&devicep->dev_splay_tree,
 436                                                  &cur_node);
 437                           if (n == NULL)
 438                             {
 439                               /* Could be possibly zero size array section.  */
 440                               cur_node.host_end--;
 441                               n = splay_tree_lookup (&devicep->dev_splay_tree,
 442                                                      &cur_node);
 443                               if (n == NULL)
 444                                 {
 445                                   cur_node.host_start--;
 446                                   n = splay_tree_lookup
 447                                         (&devicep->dev_splay_tree, &cur_node);
 448                                   cur_node.host_start++;
 449                                 }
 450                             }
 451                           if (n == NULL)
 452                             gomp_fatal ("Pointer target of array section "
 453                                         "wasn't mapped");
 454                           cur_node.host_start -= n->host_start;
 455                           cur_node.tgt_offset = n->tgt->tgt_start
 456                                                 + n->tgt_offset
 457                                                 + cur_node.host_start;
 458                           /* At this point tgt_offset is target address of the
 459                              array section.  Now subtract bias to get what we
 460                              want to initialize the pointer with.  */
 461                           cur_node.tgt_offset -= sizes[j];
 462                           devicep->host2dev_func (devicep->target_id,
 463                              (void *) (tgt->tgt_start + k->tgt_offset
 464                                        + ((uintptr_t) hostaddrs[j]
 465                                           - k->host_start)),
 466                              (void *) &cur_node.tgt_offset,
 467                              sizeof (void *));
 468                           i++;
 469                         }
 470                       break;
 471                   }
 472                 array++;
 473               }
 474           }
 475     }
 476   if (is_target)
 477     {
 478       for (i = 0; i < mapnum; i++)
 479         {
 480           if (tgt->list[i] == NULL)
 481             cur_node.tgt_offset = (uintptr_t) NULL;
 482           else
 483             cur_node.tgt_offset = tgt->list[i]->tgt->tgt_start
 484                                   + tgt->list[i]->tgt_offset;
 485           devicep->host2dev_func (devicep->target_id,
 486                                   (void *) (tgt->tgt_start
 487                                             + i * sizeof (void *)),
 488                                   (void *) &cur_node.tgt_offset,
 489                                   sizeof (void *));
 490         }
 491     }
 492
 493   gomp_mutex_unlock (&devicep->dev_env_lock);
 494   return tgt;
 495 }
 496
 497 static void
 498 gomp_unmap_tgt (struct target_mem_desc *tgt)
 499 {
 500   /* Deallocate on target the tgt->tgt_start .. tgt->tgt_end region.  */
 501   if (tgt->tgt_end)
 502     tgt->device_descr->free_func (tgt->device_descr->target_id, tgt->to_free);
 503
 504   free (tgt->array);
 505   free (tgt);
 506 }
 507
 508 static void
 509 gomp_unmap_vars (struct target_mem_desc *tgt)
 510 {
 511   struct gomp_device_descr *devicep = tgt->device_descr;
 512
 513   if (tgt->list_count == 0)
 514     {
 515       free (tgt);
 516       return;
 517     }
 518
 519   size_t i;
 520   gomp_mutex_lock (&devicep->dev_env_lock);
 521   for (i = 0; i < tgt->list_count; i++)
 522     if (tgt->list[i] == NULL)
 523       ;
 524     else if (tgt->list[i]->refcount > 1)
 525       tgt->list[i]->refcount--;
 526     else
 527       {
 528         splay_tree_key k = tgt->list[i];
 529         if (k->copy_from)
 530           devicep->dev2host_func (devicep->target_id, (void *) k->host_start,
 531                                   (void *) (k->tgt->tgt_start + k->tgt_offset),
 532                                   k->host_end - k->host_start);
 533         splay_tree_remove (&devicep->dev_splay_tree, k);
 534         if (k->tgt->refcount > 1)
 535           k->tgt->refcount--;
 536         else
 537           gomp_unmap_tgt (k->tgt);
 538       }
 539
 540   if (tgt->refcount > 1)
 541     tgt->refcount--;
 542   else
 543     gomp_unmap_tgt (tgt);
 544   gomp_mutex_unlock (&devicep->dev_env_lock);
 545 }
 546
 547 static void
 548 gomp_update (struct gomp_device_descr *devicep, size_t mapnum,
 549              void **hostaddrs, size_t *sizes, unsigned char *kinds)
 550 {
 551   size_t i;
 552   struct splay_tree_key_s cur_node;
 553
 554   if (!devicep)
 555     return;
 556
 557   if (mapnum == 0)
 558     return;
 559
 560   gomp_mutex_lock (&devicep->dev_env_lock);
 561   for (i = 0; i < mapnum; i++)
 562     if (sizes[i])
 563       {
 564         cur_node.host_start = (uintptr_t) hostaddrs[i];
 565         cur_node.host_end = cur_node.host_start + sizes[i];
 566         splay_tree_key n = splay_tree_lookup (&devicep->dev_splay_tree,
 567                                               &cur_node);
 568         if (n)
 569           {
 570             if (n->host_start > cur_node.host_start
 571                 || n->host_end < cur_node.host_end)
 572               gomp_fatal ("Trying to update [%p..%p) object when"
 573                           "only [%p..%p) is mapped",
 574                           (void *) cur_node.host_start,
 575                           (void *) cur_node.host_end,
 576                           (void *) n->host_start,
 577                           (void *) n->host_end);
 578             if ((kinds[i] & 7) == 1)
 579               devicep->host2dev_func (devicep->target_id,
 580                                       (void *) (n->tgt->tgt_start
 581                                                 + n->tgt_offset
 582                                                 + cur_node.host_start
 583                                                 - n->host_start),
 584                                       (void *) cur_node.host_start,
 585                                       cur_node.host_end - cur_node.host_start);
 586             else if ((kinds[i] & 7) == 2)
 587               devicep->dev2host_func (devicep->target_id,
 588                                       (void *) cur_node.host_start,
 589                                       (void *) (n->tgt->tgt_start
 590                                                 + n->tgt_offset
 591                                                 + cur_node.host_start
 592                                                 - n->host_start),
 593                                       cur_node.host_end - cur_node.host_start);
 594           }
 595         else
 596           gomp_fatal ("Trying to update [%p..%p) object that is not mapped",
 597                       (void *) cur_node.host_start,
 598                       (void *) cur_node.host_end);
 599       }
 600   gomp_mutex_unlock (&devicep->dev_env_lock);
 601 }
 602
 603 /* This function should be called from every offload image.
 604    It gets the descriptor of the host func and var tables HOST_TABLE, TYPE of
 605    the target, and TARGET_DATA needed by target plugin.  */
 606
 607 void
 608 GOMP_offload_register (void *host_table, enum offload_target_type target_type,
 609                        void *target_data)
 610 {
 611   offload_images = gomp_realloc (offload_images,
 612                                  (num_offload_images + 1)
 613                                  * sizeof (struct offload_image_descr));
 614
 615   offload_images[num_offload_images].type = target_type;
 616   offload_images[num_offload_images].host_table = host_table;
 617   offload_images[num_offload_images].target_data = target_data;
 618
 619   num_offload_images++;
 620 }
 621
 622 /* This function initializes the target device, specified by DEVICEP.  */
 623
 624 static void
 625 gomp_init_device (struct gomp_device_descr *devicep)
 626 {
 627   devicep->init_device_func (devicep->target_id);
 628
 629   /* Get address mapping table for device.  */
 630   struct mapping_table *table = NULL;
 631   int num_entries = devicep->get_table_func (devicep->target_id, &table);
 632
 633   /* Insert host-target address mapping into dev_splay_tree.  */
 634   int i;
 635   for (i = 0; i < num_entries; i++)
 636     {
 637       struct target_mem_desc *tgt = gomp_malloc (sizeof (*tgt));
 638       tgt->refcount = 1;
 639       tgt->array = gomp_malloc (sizeof (*tgt->array));
 640       tgt->tgt_start = table[i].tgt_start;
 641       tgt->tgt_end = table[i].tgt_end;
 642       tgt->to_free = NULL;
 643       tgt->list_count = 0;
 644       tgt->device_descr = devicep;
 645       splay_tree_node node = tgt->array;
 646       splay_tree_key k = &node->key;
 647       k->host_start = table[i].host_start;
 648       k->host_end = table[i].host_end;
 649       k->tgt_offset = 0;
 650       k->refcount = 1;
 651       k->copy_from = false;
 652       k->tgt = tgt;
 653       node->left = NULL;
 654       node->right = NULL;
 655       splay_tree_insert (&devicep->dev_splay_tree, node);
 656     }
 657
 658   free (table);
 659   devicep->is_initialized = true;
 660 }
 661
 662 /* Called when encountering a target directive.  If DEVICE
 663    is -1, it means use device-var ICV.  If it is -2 (or any other value
 664    larger than last available hw device, use host fallback.
 665    FN is address of host code, OPENMP_TARGET contains value of the
 666    __OPENMP_TARGET__ symbol in the shared library or binary that invokes
 667    GOMP_target.  HOSTADDRS, SIZES and KINDS are arrays
 668    with MAPNUM entries, with addresses of the host objects,
 669    sizes of the host objects (resp. for pointer kind pointer bias
 670    and assumed sizeof (void *) size) and kinds.  */
 671
 672 void
 673 GOMP_target (int device, void (*fn) (void *), const void *openmp_target,
 674              size_t mapnum, void **hostaddrs, size_t *sizes,
 675              unsigned char *kinds)
 676 {
 677   struct gomp_device_descr *devicep = resolve_device (device);
 678   if (devicep == NULL)
 679     {
 680       /* Host fallback.  */
 681       struct gomp_thread old_thr, *thr = gomp_thread ();
 682       old_thr = *thr;
 683       memset (thr, '\0', sizeof (*thr));
 684       if (gomp_places_list)
 685         {
 686           thr->place = old_thr.place;
 687           thr->ts.place_partition_len = gomp_places_list_len;
 688         }
 689       fn (hostaddrs);
 690       gomp_free_thread (thr);
 691       *thr = old_thr;
 692       return;
 693     }
 694
 695   gomp_mutex_lock (&devicep->dev_env_lock);
 696   if (!devicep->is_initialized)
 697     gomp_init_device (devicep);
 698
 699   struct splay_tree_key_s k;
 700   k.host_start = (uintptr_t) fn;
 701   k.host_end = k.host_start + 1;
 702   splay_tree_key tgt_fn = splay_tree_lookup (&devicep->dev_splay_tree, &k);
 703   if (tgt_fn == NULL)
 704     gomp_fatal ("Target function wasn't mapped");
 705   gomp_mutex_unlock (&devicep->dev_env_lock);
 706
 707   struct target_mem_desc *tgt_vars
 708     = gomp_map_vars (devicep, mapnum, hostaddrs, sizes, kinds, true);
 709   struct gomp_thread old_thr, *thr = gomp_thread ();
 710   old_thr = *thr;
 711   memset (thr, '\0', sizeof (*thr));
 712   if (gomp_places_list)
 713     {
 714       thr->place = old_thr.place;
 715       thr->ts.place_partition_len = gomp_places_list_len;
 716     }
 717   devicep->run_func (devicep->target_id, (void *) tgt_fn->tgt->tgt_start,
 718                      (void *) tgt_vars->tgt_start);
 719   gomp_free_thread (thr);
 720   *thr = old_thr;
 721   gomp_unmap_vars (tgt_vars);
 722 }
 723
 724 void
 725 GOMP_target_data (int device, const void *openmp_target, size_t mapnum,
 726                   void **hostaddrs, size_t *sizes, unsigned char *kinds)
 727 {
 728   struct gomp_device_descr *devicep = resolve_device (device);
 729   if (devicep == NULL)
 730     {
 731       /* Host fallback.  */
 732       struct gomp_task_icv *icv = gomp_icv (false);
 733       if (icv->target_data)
 734         {
 735           /* Even when doing a host fallback, if there are any active
 736              #pragma omp target data constructs, need to remember the
 737              new #pragma omp target data, otherwise GOMP_target_end_data
 738              would get out of sync.  */
 739           struct target_mem_desc *tgt
 740             = gomp_map_vars (NULL, 0, NULL, NULL, NULL, false);
 741           tgt->prev = icv->target_data;
 742           icv->target_data = tgt;
 743         }
 744       return;
 745     }
 746
 747   gomp_mutex_lock (&devicep->dev_env_lock);
 748   if (!devicep->is_initialized)
 749     gomp_init_device (devicep);
 750   gomp_mutex_unlock (&devicep->dev_env_lock);
 751
 752   struct target_mem_desc *tgt
 753     = gomp_map_vars (devicep, mapnum, hostaddrs, sizes, kinds, false);
 754   struct gomp_task_icv *icv = gomp_icv (true);
 755   tgt->prev = icv->target_data;
 756   icv->target_data = tgt;
 757 }
 758
 759 void
 760 GOMP_target_end_data (void)
 761 {
 762   struct gomp_task_icv *icv = gomp_icv (false);
 763   if (icv->target_data)
 764     {
 765       struct target_mem_desc *tgt = icv->target_data;
 766       icv->target_data = tgt->prev;
 767       gomp_unmap_vars (tgt);
 768     }
 769 }
 770
 771 void
 772 GOMP_target_update (int device, const void *openmp_target, size_t mapnum,
 773                     void **hostaddrs, size_t *sizes, unsigned char *kinds)
 774 {
 775   struct gomp_device_descr *devicep = resolve_device (device);
 776   if (devicep == NULL)
 777     return;
 778
 779   gomp_mutex_lock (&devicep->dev_env_lock);
 780   if (!devicep->is_initialized)
 781     gomp_init_device (devicep);
 782   gomp_mutex_unlock (&devicep->dev_env_lock);
 783
 784   gomp_update (devicep, mapnum, hostaddrs, sizes, kinds);
 785 }
 786
 787 void
 788 GOMP_teams (unsigned int num_teams, unsigned int thread_limit)
 789 {
 790   if (thread_limit)
 791     {
 792       struct gomp_task_icv *icv = gomp_icv (true);
 793       icv->thread_limit_var
 794         = thread_limit > INT_MAX ? UINT_MAX : thread_limit;
 795     }
 796   (void) num_teams;
 797 }
 798
 799 #ifdef PLUGIN_SUPPORT
 800
 801 /* This function tries to load a plugin for DEVICE.  Name of plugin is passed
 802    in PLUGIN_NAME.
 803    The handles of the found functions are stored in the corresponding fields
 804    of DEVICE.  The function returns TRUE on success and FALSE otherwise.  */
 805
 806 static bool
 807 gomp_load_plugin_for_device (struct gomp_device_descr *device,
 808                              const char *plugin_name)
 809 {
 810   void *plugin_handle = dlopen (plugin_name, RTLD_LAZY);
 811   if (!plugin_handle)
 812     return false;
 813
 814   /* Check if all required functions are available in the plugin and store
 815      their handlers.  */
 816 #define DLSYM(f)                                                    \
 817   do                                                                \
 818     {                                                               \
 819       device->f##_func = dlsym (plugin_handle, "GOMP_OFFLOAD_"#f);  \
 820       if (!device->f##_func)                                        \
 821         return false;                                               \
 822     }                                                               \
 823   while (0)
 824   DLSYM (get_type);
 825   DLSYM (get_num_devices);
 826   DLSYM (register_image);
 827   DLSYM (init_device);
 828   DLSYM (get_table);
 829   DLSYM (alloc);
 830   DLSYM (free);
 831   DLSYM (dev2host);
 832   DLSYM (host2dev);
 833   DLSYM (run);
 834 #undef DLSYM
 835
 836   return true;
 837 }
 838
 839 /* This function finds OFFLOAD_IMAGES corresponding to DEVICE type, and
 840    registers them in the plugin.  */
 841
 842 static void
 843 gomp_register_images_for_device (struct gomp_device_descr *device)
 844 {
 845   int i;
 846   for (i = 0; i < num_offload_images; i++)
 847     {
 848       struct offload_image_descr *image = &offload_images[i];
 849       if (image->type == device->type)
 850         device->register_image_func (image->host_table, image->target_data);
 851     }
 852 }
 853
 854 /* This function initializes the runtime needed for offloading.
 855    It parses the list of offload targets and tries to load the plugins for these
 856    targets.  Result of the function is properly initialized variable NUM_DEVICES
 857    and array DEVICES, containing descriptors for corresponding devices.  */
 858
 859 static void
 860 gomp_target_init (void)
 861 {
 862   const char *prefix ="libgomp-plugin-";
 863   const char *suffix = ".so.1";
 864   const char *cur, *next;
 865   char *plugin_name;
 866   int i, new_num_devices;
 867
 868   num_devices = 0;
 869   devices = NULL;
 870
 871   cur = OFFLOAD_TARGETS;
 872   if (*cur)
 873     do
 874       {
 875         struct gomp_device_descr current_device;
 876
 877         next = strchr (cur, ',');
 878
 879         plugin_name = (char *) malloc (1 + (next ? next - cur : strlen (cur))
 880                                        + strlen (prefix) + strlen (suffix));
 881         if (!plugin_name)
 882           {
 883             num_devices = 0;
 884             break;
 885           }
 886
 887         strcpy (plugin_name, prefix);
 888         strncat (plugin_name, cur, next ? next - cur : strlen (cur));
 889         strcat (plugin_name, suffix);
 890
 891         if (gomp_load_plugin_for_device (&current_device, plugin_name))
 892           {
 893             new_num_devices = current_device.get_num_devices_func ();
 894             if (new_num_devices >= 1)
 895               {
 896                 devices = realloc (devices, (num_devices + new_num_devices)
 897                                    * sizeof (struct gomp_device_descr));
 898                 if (!devices)
 899                   {
 900                     num_devices = 0;
 901                     free (plugin_name);
 902                     break;
 903                   }
 904
 905                 current_device.type = current_device.get_type_func ();
 906                 current_device.is_initialized = false;
 907                 current_device.dev_splay_tree.root = NULL;
 908                 gomp_register_images_for_device (&current_device);
 909                 for (i = 0; i < new_num_devices; i++)
 910                   {
 911                     current_device.id = num_devices + 1;
 912                     current_device.target_id = i;
 913                     devices[num_devices] = current_device;
 914                     gomp_mutex_init (&devices[num_devices].dev_env_lock);
 915                     num_devices++;
 916                   }
 917               }
 918           }
 919
 920         free (plugin_name);
 921         cur = next + 1;
 922       }
 923     while (next);
 924
 925   free (offload_images);
 926   offload_images = NULL;
 927   num_offload_images = 0;
 928 }
 929
 930 #else /* PLUGIN_SUPPORT */
 931 /* If dlfcn.h is unavailable we always fallback to host execution.
 932    GOMP_target* routines are just stubs for this case.  */
 933 static void
 934 gomp_target_init (void)
 935 {
 936 }
 937 #endif /* PLUGIN_SUPPORT */