1 /* Copyright (C) 2013-2015 Free Software Foundation, Inc.
2 Contributed by Jakub Jelinek <jakub@redhat.com>.
4 This file is part of the GNU Offloading and Multi Processing Library
7 Libgomp is free software; you can redistribute it and/or modify it
8 under the terms of the GNU General Public License as published by
9 the Free Software Foundation; either version 3, or (at your option)
12 Libgomp is distributed in the hope that it will be useful, but WITHOUT ANY
13 WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS
14 FOR A PARTICULAR PURPOSE. See the GNU General Public License for
17 Under Section 7 of GPL version 3, you are granted additional
18 permissions described in the GCC Runtime Library Exception, version
19 3.1, as published by the Free Software Foundation.
21 You should have received a copy of the GNU General Public License and
22 a copy of the GCC Runtime Library Exception along with this program;
23 see the files COPYING3 and COPYING.RUNTIME respectively. If not, see
24 <http://www.gnu.org/licenses/>. */
26 /* This file contains the support of offloading. */
30 #include "libgomp_target.h"
40 static void gomp_target_init (void);
42 static pthread_once_t gomp_is_initialized
= PTHREAD_ONCE_INIT
;
44 /* Forward declaration for a node in the tree. */
45 typedef struct splay_tree_node_s
*splay_tree_node
;
46 typedef struct splay_tree_s
*splay_tree
;
47 typedef struct splay_tree_key_s
*splay_tree_key
;
49 struct target_mem_desc
{
50 /* Reference count. */
52 /* All the splay nodes allocated together. */
53 splay_tree_node array
;
54 /* Start of the target region. */
56 /* End of the targer region. */
60 /* Previous target_mem_desc. */
61 struct target_mem_desc
*prev
;
62 /* Number of items in following list. */
65 /* Corresponding target device descriptor. */
66 struct gomp_device_descr
*device_descr
;
68 /* List of splay keys to remove (or decrease refcount)
69 at the end of region. */
70 splay_tree_key list
[];
73 struct splay_tree_key_s
{
74 /* Address of the host object. */
76 /* Address immediately after the host object. */
78 /* Descriptor of the target memory. */
79 struct target_mem_desc
*tgt
;
80 /* Offset from tgt->tgt_start to the start of the target object. */
82 /* Reference count. */
84 /* True if data should be copied from device to host at the end. */
88 /* This structure describes an offload image.
89 It contains type of the target device, pointer to host table descriptor, and
90 pointer to target data. */
91 struct offload_image_descr
{
92 enum offload_target_type type
;
97 /* Array of descriptors of offload images. */
98 static struct offload_image_descr
*offload_images
;
100 /* Total number of offload images. */
101 static int num_offload_images
;
103 /* Array of descriptors for all available devices. */
104 static struct gomp_device_descr
*devices
;
106 /* Total number of available devices. */
107 static int num_devices
;
109 /* The comparison function. */
112 splay_compare (splay_tree_key x
, splay_tree_key y
)
114 if (x
->host_start
== x
->host_end
115 && y
->host_start
== y
->host_end
)
117 if (x
->host_end
<= y
->host_start
)
119 if (x
->host_start
>= y
->host_end
)
124 #include "splay-tree.h"
126 /* This structure describes accelerator device.
127 It contains ID-number of the device, its type, function handlers for
128 interaction with the device, and information about mapped memory. */
129 struct gomp_device_descr
131 /* This is the ID number of device. It could be specified in DEVICE-clause of
135 /* This is the ID number of device among devices of the same type. */
138 /* This is the TYPE of device. */
139 enum offload_target_type type
;
141 /* Set to true when device is initialized. */
144 /* Function handlers. */
145 int (*get_type_func
) (void);
146 int (*get_num_devices_func
) (void);
147 void (*register_image_func
) (void *, void *);
148 void (*init_device_func
) (int);
149 int (*get_table_func
) (int, void *);
150 void *(*alloc_func
) (int, size_t);
151 void (*free_func
) (int, void *);
152 void *(*host2dev_func
) (int, void *, const void *, size_t);
153 void *(*dev2host_func
) (int, void *, const void *, size_t);
154 void (*run_func
) (int, void *, void *);
156 /* Splay tree containing information about mapped memory regions. */
157 struct splay_tree_s dev_splay_tree
;
159 /* Mutex for operating with the splay tree and other shared structures. */
160 gomp_mutex_t dev_env_lock
;
164 gomp_get_num_devices (void)
166 (void) pthread_once (&gomp_is_initialized
, gomp_target_init
);
170 static struct gomp_device_descr
*
171 resolve_device (int device_id
)
175 struct gomp_task_icv
*icv
= gomp_icv (false);
176 device_id
= icv
->default_device_var
;
179 if (device_id
< 0 || device_id
>= gomp_get_num_devices ())
182 return &devices
[device_id
];
186 /* Handle the case where splay_tree_lookup found oldn for newn.
187 Helper function of gomp_map_vars. */
190 gomp_map_vars_existing (splay_tree_key oldn
, splay_tree_key newn
,
193 if (oldn
->host_start
> newn
->host_start
194 || oldn
->host_end
< newn
->host_end
)
195 gomp_fatal ("Trying to map into device [%p..%p) object when"
196 "[%p..%p) is already mapped",
197 (void *) newn
->host_start
, (void *) newn
->host_end
,
198 (void *) oldn
->host_start
, (void *) oldn
->host_end
);
202 static struct target_mem_desc
*
203 gomp_map_vars (struct gomp_device_descr
*devicep
, size_t mapnum
,
204 void **hostaddrs
, size_t *sizes
, unsigned char *kinds
,
207 size_t i
, tgt_align
, tgt_size
, not_found_cnt
= 0;
208 struct splay_tree_key_s cur_node
;
209 struct target_mem_desc
*tgt
210 = gomp_malloc (sizeof (*tgt
) + sizeof (tgt
->list
[0]) * mapnum
);
211 tgt
->list_count
= mapnum
;
213 tgt
->device_descr
= devicep
;
218 tgt_align
= sizeof (void *);
222 size_t align
= 4 * sizeof (void *);
224 tgt_size
= mapnum
* sizeof (void *);
227 gomp_mutex_lock (&devicep
->dev_env_lock
);
228 for (i
= 0; i
< mapnum
; i
++)
230 if (hostaddrs
[i
] == NULL
)
235 cur_node
.host_start
= (uintptr_t) hostaddrs
[i
];
236 if ((kinds
[i
] & 7) != 4)
237 cur_node
.host_end
= cur_node
.host_start
+ sizes
[i
];
239 cur_node
.host_end
= cur_node
.host_start
+ sizeof (void *);
240 splay_tree_key n
= splay_tree_lookup (&devicep
->dev_splay_tree
,
245 gomp_map_vars_existing (n
, &cur_node
, kinds
[i
]);
249 size_t align
= (size_t) 1 << (kinds
[i
] >> 3);
252 if (tgt_align
< align
)
254 tgt_size
= (tgt_size
+ align
- 1) & ~(align
- 1);
255 tgt_size
+= cur_node
.host_end
- cur_node
.host_start
;
256 if ((kinds
[i
] & 7) == 5)
259 for (j
= i
+ 1; j
< mapnum
; j
++)
260 if ((kinds
[j
] & 7) != 4)
262 else if ((uintptr_t) hostaddrs
[j
] < cur_node
.host_start
263 || ((uintptr_t) hostaddrs
[j
] + sizeof (void *)
264 > cur_node
.host_end
))
275 if (not_found_cnt
|| is_target
)
277 /* Allocate tgt_align aligned tgt_size block of memory. */
278 /* FIXME: Perhaps change interface to allocate properly aligned
280 tgt
->to_free
= devicep
->alloc_func (devicep
->target_id
,
281 tgt_size
+ tgt_align
- 1);
282 tgt
->tgt_start
= (uintptr_t) tgt
->to_free
;
283 tgt
->tgt_start
= (tgt
->tgt_start
+ tgt_align
- 1) & ~(tgt_align
- 1);
284 tgt
->tgt_end
= tgt
->tgt_start
+ tgt_size
;
295 tgt_size
= mapnum
* sizeof (void *);
300 tgt
->array
= gomp_malloc (not_found_cnt
* sizeof (*tgt
->array
));
301 splay_tree_node array
= tgt
->array
;
304 for (i
= 0; i
< mapnum
; i
++)
305 if (tgt
->list
[i
] == NULL
)
307 if (hostaddrs
[i
] == NULL
)
309 splay_tree_key k
= &array
->key
;
310 k
->host_start
= (uintptr_t) hostaddrs
[i
];
311 if ((kinds
[i
] & 7) != 4)
312 k
->host_end
= k
->host_start
+ sizes
[i
];
314 k
->host_end
= k
->host_start
+ sizeof (void *);
316 = splay_tree_lookup (&devicep
->dev_splay_tree
, k
);
320 gomp_map_vars_existing (n
, k
, kinds
[i
]);
324 size_t align
= (size_t) 1 << (kinds
[i
] >> 3);
326 tgt_size
= (tgt_size
+ align
- 1) & ~(align
- 1);
328 k
->tgt_offset
= tgt_size
;
329 tgt_size
+= k
->host_end
- k
->host_start
;
330 k
->copy_from
= false;
331 if ((kinds
[i
] & 7) == 2 || (kinds
[i
] & 7) == 3)
337 splay_tree_insert (&devicep
->dev_splay_tree
, array
);
338 switch (kinds
[i
] & 7)
345 /* FIXME: Perhaps add some smarts, like if copying
346 several adjacent fields from host to target, use some
347 host buffer to avoid sending each var individually. */
348 devicep
->host2dev_func (devicep
->target_id
,
349 (void *) (tgt
->tgt_start
351 (void *) k
->host_start
,
352 k
->host_end
- k
->host_start
);
354 case 4: /* POINTER */
356 = (uintptr_t) *(void **) k
->host_start
;
357 if (cur_node
.host_start
== (uintptr_t) NULL
)
359 cur_node
.tgt_offset
= (uintptr_t) NULL
;
360 devicep
->host2dev_func (devicep
->target_id
,
361 (void *) (tgt
->tgt_start
363 (void *) &cur_node
.tgt_offset
,
367 /* Add bias to the pointer value. */
368 cur_node
.host_start
+= sizes
[i
];
369 cur_node
.host_end
= cur_node
.host_start
+ 1;
370 n
= splay_tree_lookup (&devicep
->dev_splay_tree
,
374 /* Could be possibly zero size array section. */
376 n
= splay_tree_lookup (&devicep
->dev_splay_tree
,
380 cur_node
.host_start
--;
381 n
= splay_tree_lookup (&devicep
->dev_splay_tree
,
383 cur_node
.host_start
++;
387 gomp_fatal ("Pointer target of array section "
389 cur_node
.host_start
-= n
->host_start
;
390 cur_node
.tgt_offset
= n
->tgt
->tgt_start
+ n
->tgt_offset
391 + cur_node
.host_start
;
392 /* At this point tgt_offset is target address of the
393 array section. Now subtract bias to get what we want
394 to initialize the pointer with. */
395 cur_node
.tgt_offset
-= sizes
[i
];
396 devicep
->host2dev_func (devicep
->target_id
,
397 (void *) (tgt
->tgt_start
399 (void *) &cur_node
.tgt_offset
,
402 case 5: /* TO_PSET */
403 devicep
->host2dev_func (devicep
->target_id
,
404 (void *) (tgt
->tgt_start
406 (void *) k
->host_start
,
407 k
->host_end
- k
->host_start
);
408 for (j
= i
+ 1; j
< mapnum
; j
++)
409 if ((kinds
[j
] & 7) != 4)
411 else if ((uintptr_t) hostaddrs
[j
] < k
->host_start
412 || ((uintptr_t) hostaddrs
[j
] + sizeof (void *)
420 = (uintptr_t) *(void **) hostaddrs
[j
];
421 if (cur_node
.host_start
== (uintptr_t) NULL
)
423 cur_node
.tgt_offset
= (uintptr_t) NULL
;
424 devicep
->host2dev_func (devicep
->target_id
,
425 (void *) (tgt
->tgt_start
+ k
->tgt_offset
426 + ((uintptr_t) hostaddrs
[j
]
428 (void *) &cur_node
.tgt_offset
,
433 /* Add bias to the pointer value. */
434 cur_node
.host_start
+= sizes
[j
];
435 cur_node
.host_end
= cur_node
.host_start
+ 1;
436 n
= splay_tree_lookup (&devicep
->dev_splay_tree
,
440 /* Could be possibly zero size array section. */
442 n
= splay_tree_lookup (&devicep
->dev_splay_tree
,
446 cur_node
.host_start
--;
447 n
= splay_tree_lookup
448 (&devicep
->dev_splay_tree
, &cur_node
);
449 cur_node
.host_start
++;
453 gomp_fatal ("Pointer target of array section "
455 cur_node
.host_start
-= n
->host_start
;
456 cur_node
.tgt_offset
= n
->tgt
->tgt_start
458 + cur_node
.host_start
;
459 /* At this point tgt_offset is target address of the
460 array section. Now subtract bias to get what we
461 want to initialize the pointer with. */
462 cur_node
.tgt_offset
-= sizes
[j
];
463 devicep
->host2dev_func (devicep
->target_id
,
464 (void *) (tgt
->tgt_start
+ k
->tgt_offset
465 + ((uintptr_t) hostaddrs
[j
]
467 (void *) &cur_node
.tgt_offset
,
479 for (i
= 0; i
< mapnum
; i
++)
481 if (tgt
->list
[i
] == NULL
)
482 cur_node
.tgt_offset
= (uintptr_t) NULL
;
484 cur_node
.tgt_offset
= tgt
->list
[i
]->tgt
->tgt_start
485 + tgt
->list
[i
]->tgt_offset
;
486 devicep
->host2dev_func (devicep
->target_id
,
487 (void *) (tgt
->tgt_start
488 + i
* sizeof (void *)),
489 (void *) &cur_node
.tgt_offset
,
494 gomp_mutex_unlock (&devicep
->dev_env_lock
);
499 gomp_unmap_tgt (struct target_mem_desc
*tgt
)
501 /* Deallocate on target the tgt->tgt_start .. tgt->tgt_end region. */
503 tgt
->device_descr
->free_func (tgt
->device_descr
->target_id
, tgt
->to_free
);
510 gomp_unmap_vars (struct target_mem_desc
*tgt
)
512 struct gomp_device_descr
*devicep
= tgt
->device_descr
;
514 if (tgt
->list_count
== 0)
521 gomp_mutex_lock (&devicep
->dev_env_lock
);
522 for (i
= 0; i
< tgt
->list_count
; i
++)
523 if (tgt
->list
[i
] == NULL
)
525 else if (tgt
->list
[i
]->refcount
> 1)
526 tgt
->list
[i
]->refcount
--;
529 splay_tree_key k
= tgt
->list
[i
];
531 devicep
->dev2host_func (devicep
->target_id
, (void *) k
->host_start
,
532 (void *) (k
->tgt
->tgt_start
+ k
->tgt_offset
),
533 k
->host_end
- k
->host_start
);
534 splay_tree_remove (&devicep
->dev_splay_tree
, k
);
535 if (k
->tgt
->refcount
> 1)
538 gomp_unmap_tgt (k
->tgt
);
541 if (tgt
->refcount
> 1)
544 gomp_unmap_tgt (tgt
);
545 gomp_mutex_unlock (&devicep
->dev_env_lock
);
549 gomp_update (struct gomp_device_descr
*devicep
, size_t mapnum
,
550 void **hostaddrs
, size_t *sizes
, unsigned char *kinds
)
553 struct splay_tree_key_s cur_node
;
561 gomp_mutex_lock (&devicep
->dev_env_lock
);
562 for (i
= 0; i
< mapnum
; i
++)
565 cur_node
.host_start
= (uintptr_t) hostaddrs
[i
];
566 cur_node
.host_end
= cur_node
.host_start
+ sizes
[i
];
567 splay_tree_key n
= splay_tree_lookup (&devicep
->dev_splay_tree
,
571 if (n
->host_start
> cur_node
.host_start
572 || n
->host_end
< cur_node
.host_end
)
573 gomp_fatal ("Trying to update [%p..%p) object when"
574 "only [%p..%p) is mapped",
575 (void *) cur_node
.host_start
,
576 (void *) cur_node
.host_end
,
577 (void *) n
->host_start
,
578 (void *) n
->host_end
);
579 if ((kinds
[i
] & 7) == 1)
580 devicep
->host2dev_func (devicep
->target_id
,
581 (void *) (n
->tgt
->tgt_start
583 + cur_node
.host_start
585 (void *) cur_node
.host_start
,
586 cur_node
.host_end
- cur_node
.host_start
);
587 else if ((kinds
[i
] & 7) == 2)
588 devicep
->dev2host_func (devicep
->target_id
,
589 (void *) cur_node
.host_start
,
590 (void *) (n
->tgt
->tgt_start
592 + cur_node
.host_start
594 cur_node
.host_end
- cur_node
.host_start
);
597 gomp_fatal ("Trying to update [%p..%p) object that is not mapped",
598 (void *) cur_node
.host_start
,
599 (void *) cur_node
.host_end
);
601 gomp_mutex_unlock (&devicep
->dev_env_lock
);
604 /* This function should be called from every offload image.
605 It gets the descriptor of the host func and var tables HOST_TABLE, TYPE of
606 the target, and TARGET_DATA needed by target plugin. */
609 GOMP_offload_register (void *host_table
, enum offload_target_type target_type
,
612 offload_images
= gomp_realloc (offload_images
,
613 (num_offload_images
+ 1)
614 * sizeof (struct offload_image_descr
));
616 offload_images
[num_offload_images
].type
= target_type
;
617 offload_images
[num_offload_images
].host_table
= host_table
;
618 offload_images
[num_offload_images
].target_data
= target_data
;
620 num_offload_images
++;
623 /* This function initializes the target device, specified by DEVICEP. */
626 gomp_init_device (struct gomp_device_descr
*devicep
)
628 devicep
->init_device_func (devicep
->target_id
);
630 /* Get address mapping table for device. */
631 struct mapping_table
*table
= NULL
;
632 int num_entries
= devicep
->get_table_func (devicep
->target_id
, &table
);
634 /* Insert host-target address mapping into dev_splay_tree. */
636 for (i
= 0; i
< num_entries
; i
++)
638 struct target_mem_desc
*tgt
= gomp_malloc (sizeof (*tgt
));
640 tgt
->array
= gomp_malloc (sizeof (*tgt
->array
));
641 tgt
->tgt_start
= table
[i
].tgt_start
;
642 tgt
->tgt_end
= table
[i
].tgt_end
;
645 tgt
->device_descr
= devicep
;
646 splay_tree_node node
= tgt
->array
;
647 splay_tree_key k
= &node
->key
;
648 k
->host_start
= table
[i
].host_start
;
649 k
->host_end
= table
[i
].host_end
;
652 k
->copy_from
= false;
656 splay_tree_insert (&devicep
->dev_splay_tree
, node
);
660 devicep
->is_initialized
= true;
663 /* Called when encountering a target directive. If DEVICE
664 is -1, it means use device-var ICV. If it is -2 (or any other value
665 larger than last available hw device, use host fallback.
666 FN is address of host code, OPENMP_TARGET contains value of the
667 __OPENMP_TARGET__ symbol in the shared library or binary that invokes
668 GOMP_target. HOSTADDRS, SIZES and KINDS are arrays
669 with MAPNUM entries, with addresses of the host objects,
670 sizes of the host objects (resp. for pointer kind pointer bias
671 and assumed sizeof (void *) size) and kinds. */
674 GOMP_target (int device
, void (*fn
) (void *), const void *openmp_target
,
675 size_t mapnum
, void **hostaddrs
, size_t *sizes
,
676 unsigned char *kinds
)
678 struct gomp_device_descr
*devicep
= resolve_device (device
);
682 struct gomp_thread old_thr
, *thr
= gomp_thread ();
684 memset (thr
, '\0', sizeof (*thr
));
685 if (gomp_places_list
)
687 thr
->place
= old_thr
.place
;
688 thr
->ts
.place_partition_len
= gomp_places_list_len
;
691 gomp_free_thread (thr
);
696 gomp_mutex_lock (&devicep
->dev_env_lock
);
697 if (!devicep
->is_initialized
)
698 gomp_init_device (devicep
);
700 struct splay_tree_key_s k
;
701 k
.host_start
= (uintptr_t) fn
;
702 k
.host_end
= k
.host_start
+ 1;
703 splay_tree_key tgt_fn
= splay_tree_lookup (&devicep
->dev_splay_tree
, &k
);
705 gomp_fatal ("Target function wasn't mapped");
706 gomp_mutex_unlock (&devicep
->dev_env_lock
);
708 struct target_mem_desc
*tgt_vars
709 = gomp_map_vars (devicep
, mapnum
, hostaddrs
, sizes
, kinds
, true);
710 struct gomp_thread old_thr
, *thr
= gomp_thread ();
712 memset (thr
, '\0', sizeof (*thr
));
713 if (gomp_places_list
)
715 thr
->place
= old_thr
.place
;
716 thr
->ts
.place_partition_len
= gomp_places_list_len
;
718 devicep
->run_func (devicep
->target_id
, (void *) tgt_fn
->tgt
->tgt_start
,
719 (void *) tgt_vars
->tgt_start
);
720 gomp_free_thread (thr
);
722 gomp_unmap_vars (tgt_vars
);
726 GOMP_target_data (int device
, const void *openmp_target
, size_t mapnum
,
727 void **hostaddrs
, size_t *sizes
, unsigned char *kinds
)
729 struct gomp_device_descr
*devicep
= resolve_device (device
);
733 struct gomp_task_icv
*icv
= gomp_icv (false);
734 if (icv
->target_data
)
736 /* Even when doing a host fallback, if there are any active
737 #pragma omp target data constructs, need to remember the
738 new #pragma omp target data, otherwise GOMP_target_end_data
739 would get out of sync. */
740 struct target_mem_desc
*tgt
741 = gomp_map_vars (NULL
, 0, NULL
, NULL
, NULL
, false);
742 tgt
->prev
= icv
->target_data
;
743 icv
->target_data
= tgt
;
748 gomp_mutex_lock (&devicep
->dev_env_lock
);
749 if (!devicep
->is_initialized
)
750 gomp_init_device (devicep
);
751 gomp_mutex_unlock (&devicep
->dev_env_lock
);
753 struct target_mem_desc
*tgt
754 = gomp_map_vars (devicep
, mapnum
, hostaddrs
, sizes
, kinds
, false);
755 struct gomp_task_icv
*icv
= gomp_icv (true);
756 tgt
->prev
= icv
->target_data
;
757 icv
->target_data
= tgt
;
761 GOMP_target_end_data (void)
763 struct gomp_task_icv
*icv
= gomp_icv (false);
764 if (icv
->target_data
)
766 struct target_mem_desc
*tgt
= icv
->target_data
;
767 icv
->target_data
= tgt
->prev
;
768 gomp_unmap_vars (tgt
);
773 GOMP_target_update (int device
, const void *openmp_target
, size_t mapnum
,
774 void **hostaddrs
, size_t *sizes
, unsigned char *kinds
)
776 struct gomp_device_descr
*devicep
= resolve_device (device
);
780 gomp_mutex_lock (&devicep
->dev_env_lock
);
781 if (!devicep
->is_initialized
)
782 gomp_init_device (devicep
);
783 gomp_mutex_unlock (&devicep
->dev_env_lock
);
785 gomp_update (devicep
, mapnum
, hostaddrs
, sizes
, kinds
);
789 GOMP_teams (unsigned int num_teams
, unsigned int thread_limit
)
793 struct gomp_task_icv
*icv
= gomp_icv (true);
794 icv
->thread_limit_var
795 = thread_limit
> INT_MAX
? UINT_MAX
: thread_limit
;
800 #ifdef PLUGIN_SUPPORT
802 /* This function tries to load a plugin for DEVICE. Name of plugin is passed
804 The handles of the found functions are stored in the corresponding fields
805 of DEVICE. The function returns TRUE on success and FALSE otherwise. */
808 gomp_load_plugin_for_device (struct gomp_device_descr
*device
,
809 const char *plugin_name
)
811 void *plugin_handle
= dlopen (plugin_name
, RTLD_LAZY
);
815 /* Check if all required functions are available in the plugin and store
820 device->f##_func = dlsym (plugin_handle, "GOMP_OFFLOAD_"#f); \
821 if (!device->f##_func) \
826 DLSYM (get_num_devices
);
827 DLSYM (register_image
);
840 /* This function finds OFFLOAD_IMAGES corresponding to DEVICE type, and
841 registers them in the plugin. */
844 gomp_register_images_for_device (struct gomp_device_descr
*device
)
847 for (i
= 0; i
< num_offload_images
; i
++)
849 struct offload_image_descr
*image
= &offload_images
[i
];
850 if (image
->type
== device
->type
)
851 device
->register_image_func (image
->host_table
, image
->target_data
);
855 /* This function initializes the runtime needed for offloading.
856 It parses the list of offload targets and tries to load the plugins for these
857 targets. Result of the function is properly initialized variable NUM_DEVICES
858 and array DEVICES, containing descriptors for corresponding devices. */
861 gomp_target_init (void)
863 const char *prefix
="libgomp-plugin-";
864 const char *suffix
= ".so.1";
865 const char *cur
, *next
;
867 int i
, new_num_devices
;
872 cur
= OFFLOAD_TARGETS
;
876 struct gomp_device_descr current_device
;
878 next
= strchr (cur
, ',');
880 plugin_name
= (char *) malloc (1 + (next
? next
- cur
: strlen (cur
))
881 + strlen (prefix
) + strlen (suffix
));
888 strcpy (plugin_name
, prefix
);
889 strncat (plugin_name
, cur
, next
? next
- cur
: strlen (cur
));
890 strcat (plugin_name
, suffix
);
892 if (gomp_load_plugin_for_device (¤t_device
, plugin_name
))
894 new_num_devices
= current_device
.get_num_devices_func ();
895 if (new_num_devices
>= 1)
897 devices
= realloc (devices
, (num_devices
+ new_num_devices
)
898 * sizeof (struct gomp_device_descr
));
906 current_device
.type
= current_device
.get_type_func ();
907 current_device
.is_initialized
= false;
908 current_device
.dev_splay_tree
.root
= NULL
;
909 gomp_register_images_for_device (¤t_device
);
910 for (i
= 0; i
< new_num_devices
; i
++)
912 current_device
.id
= num_devices
+ 1;
913 current_device
.target_id
= i
;
914 devices
[num_devices
] = current_device
;
915 gomp_mutex_init (&devices
[num_devices
].dev_env_lock
);
926 free (offload_images
);
927 offload_images
= NULL
;
928 num_offload_images
= 0;
931 #else /* PLUGIN_SUPPORT */
932 /* If dlfcn.h is unavailable we always fallback to host execution.
933 GOMP_target* routines are just stubs for this case. */
935 gomp_target_init (void)
938 #endif /* PLUGIN_SUPPORT */