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 "oacc-plugin.h"
32 #include "gomp-constants.h"
43 static void gomp_target_init (void);
45 /* The whole initialization code for offloading plugins is only run one. */
46 static pthread_once_t gomp_is_initialized
= PTHREAD_ONCE_INIT
;
48 /* This structure describes an offload image.
49 It contains type of the target device, pointer to host table descriptor, and
50 pointer to target data. */
51 struct offload_image_descr
{
52 enum offload_target_type type
;
57 /* Array of descriptors of offload images. */
58 static struct offload_image_descr
*offload_images
;
60 /* Total number of offload images. */
61 static int num_offload_images
;
63 /* Array of descriptors for all available devices. */
64 static struct gomp_device_descr
*devices
;
67 /* Total number of available devices. */
68 static int num_devices
;
71 /* Number of GOMP_OFFLOAD_CAP_OPENMP_400 devices. */
72 static int num_devices_openmp
;
74 /* The comparison function. */
77 splay_compare (splay_tree_key x
, splay_tree_key y
)
79 if (x
->host_start
== x
->host_end
80 && y
->host_start
== y
->host_end
)
82 if (x
->host_end
<= y
->host_start
)
84 if (x
->host_start
>= y
->host_end
)
89 #include "splay-tree.h"
92 gomp_init_targets_once (void)
94 (void) pthread_once (&gomp_is_initialized
, gomp_target_init
);
98 gomp_get_num_devices (void)
100 gomp_init_targets_once ();
101 return num_devices_openmp
;
104 static struct gomp_device_descr
*
105 resolve_device (int device_id
)
107 if (device_id
== GOMP_DEVICE_ICV
)
109 struct gomp_task_icv
*icv
= gomp_icv (false);
110 device_id
= icv
->default_device_var
;
113 if (device_id
< 0 || device_id
>= gomp_get_num_devices ())
116 return &devices
[device_id
];
120 /* Handle the case where splay_tree_lookup found oldn for newn.
121 Helper function of gomp_map_vars. */
124 gomp_map_vars_existing (splay_tree_key oldn
, splay_tree_key newn
,
127 if ((kind
& GOMP_MAP_FLAG_FORCE
)
128 || oldn
->host_start
> newn
->host_start
129 || oldn
->host_end
< newn
->host_end
)
130 gomp_fatal ("Trying to map into device [%p..%p) object when "
131 "[%p..%p) is already mapped",
132 (void *) newn
->host_start
, (void *) newn
->host_end
,
133 (void *) oldn
->host_start
, (void *) oldn
->host_end
);
138 get_kind (bool is_openacc
, void *kinds
, int idx
)
140 return is_openacc
? ((unsigned short *) kinds
)[idx
]
141 : ((unsigned char *) kinds
)[idx
];
144 attribute_hidden
struct target_mem_desc
*
145 gomp_map_vars (struct gomp_device_descr
*devicep
, size_t mapnum
,
146 void **hostaddrs
, void **devaddrs
, size_t *sizes
, void *kinds
,
147 bool is_openacc
, bool is_target
)
149 size_t i
, tgt_align
, tgt_size
, not_found_cnt
= 0;
150 const int rshift
= is_openacc
? 8 : 3;
151 const int typemask
= is_openacc
? 0xff : 0x7;
152 struct gomp_memory_mapping
*mm
= &devicep
->mem_map
;
153 struct splay_tree_key_s cur_node
;
154 struct target_mem_desc
*tgt
155 = gomp_malloc (sizeof (*tgt
) + sizeof (tgt
->list
[0]) * mapnum
);
156 tgt
->list_count
= mapnum
;
158 tgt
->device_descr
= devicep
;
164 tgt_align
= sizeof (void *);
168 size_t align
= 4 * sizeof (void *);
170 tgt_size
= mapnum
* sizeof (void *);
173 gomp_mutex_lock (&mm
->lock
);
175 for (i
= 0; i
< mapnum
; i
++)
177 int kind
= get_kind (is_openacc
, kinds
, i
);
178 if (hostaddrs
[i
] == NULL
)
183 cur_node
.host_start
= (uintptr_t) hostaddrs
[i
];
184 if (!GOMP_MAP_POINTER_P (kind
& typemask
))
185 cur_node
.host_end
= cur_node
.host_start
+ sizes
[i
];
187 cur_node
.host_end
= cur_node
.host_start
+ sizeof (void *);
188 splay_tree_key n
= splay_tree_lookup (&mm
->splay_tree
, &cur_node
);
192 gomp_map_vars_existing (n
, &cur_node
, kind
& typemask
);
198 size_t align
= (size_t) 1 << (kind
>> rshift
);
200 if (tgt_align
< align
)
202 tgt_size
= (tgt_size
+ align
- 1) & ~(align
- 1);
203 tgt_size
+= cur_node
.host_end
- cur_node
.host_start
;
204 if ((kind
& typemask
) == GOMP_MAP_TO_PSET
)
207 for (j
= i
+ 1; j
< mapnum
; j
++)
208 if (!GOMP_MAP_POINTER_P (get_kind (is_openacc
, kinds
, j
)
211 else if ((uintptr_t) hostaddrs
[j
] < cur_node
.host_start
212 || ((uintptr_t) hostaddrs
[j
] + sizeof (void *)
213 > cur_node
.host_end
))
227 gomp_fatal ("unexpected aggregation");
228 tgt
->to_free
= devaddrs
[0];
229 tgt
->tgt_start
= (uintptr_t) tgt
->to_free
;
230 tgt
->tgt_end
= tgt
->tgt_start
+ sizes
[0];
232 else if (not_found_cnt
|| is_target
)
234 /* Allocate tgt_align aligned tgt_size block of memory. */
235 /* FIXME: Perhaps change interface to allocate properly aligned
237 tgt
->to_free
= devicep
->alloc_func (devicep
->target_id
,
238 tgt_size
+ tgt_align
- 1);
239 tgt
->tgt_start
= (uintptr_t) tgt
->to_free
;
240 tgt
->tgt_start
= (tgt
->tgt_start
+ tgt_align
- 1) & ~(tgt_align
- 1);
241 tgt
->tgt_end
= tgt
->tgt_start
+ tgt_size
;
252 tgt_size
= mapnum
* sizeof (void *);
257 tgt
->array
= gomp_malloc (not_found_cnt
* sizeof (*tgt
->array
));
258 splay_tree_node array
= tgt
->array
;
261 for (i
= 0; i
< mapnum
; i
++)
262 if (tgt
->list
[i
] == NULL
)
264 int kind
= get_kind (is_openacc
, kinds
, i
);
265 if (hostaddrs
[i
] == NULL
)
267 splay_tree_key k
= &array
->key
;
268 k
->host_start
= (uintptr_t) hostaddrs
[i
];
269 if (!GOMP_MAP_POINTER_P (kind
& typemask
))
270 k
->host_end
= k
->host_start
+ sizes
[i
];
272 k
->host_end
= k
->host_start
+ sizeof (void *);
273 splay_tree_key n
= splay_tree_lookup (&mm
->splay_tree
, k
);
277 gomp_map_vars_existing (n
, k
, kind
& typemask
);
281 size_t align
= (size_t) 1 << (kind
>> rshift
);
283 tgt_size
= (tgt_size
+ align
- 1) & ~(align
- 1);
285 k
->tgt_offset
= tgt_size
;
286 tgt_size
+= k
->host_end
- k
->host_start
;
287 k
->copy_from
= GOMP_MAP_COPY_FROM_P (kind
& typemask
);
289 k
->async_refcount
= 0;
293 splay_tree_insert (&mm
->splay_tree
, array
);
294 switch (kind
& typemask
)
298 case GOMP_MAP_FORCE_ALLOC
:
299 case GOMP_MAP_FORCE_FROM
:
302 case GOMP_MAP_TOFROM
:
303 case GOMP_MAP_FORCE_TO
:
304 case GOMP_MAP_FORCE_TOFROM
:
305 /* FIXME: Perhaps add some smarts, like if copying
306 several adjacent fields from host to target, use some
307 host buffer to avoid sending each var individually. */
308 devicep
->host2dev_func (devicep
->target_id
,
309 (void *) (tgt
->tgt_start
311 (void *) k
->host_start
,
312 k
->host_end
- k
->host_start
);
314 case GOMP_MAP_POINTER
:
316 = (uintptr_t) *(void **) k
->host_start
;
317 if (cur_node
.host_start
== (uintptr_t) NULL
)
319 cur_node
.tgt_offset
= (uintptr_t) NULL
;
320 /* FIXME: see above FIXME comment. */
321 devicep
->host2dev_func (devicep
->target_id
,
322 (void *) (tgt
->tgt_start
324 (void *) &cur_node
.tgt_offset
,
328 /* Add bias to the pointer value. */
329 cur_node
.host_start
+= sizes
[i
];
330 cur_node
.host_end
= cur_node
.host_start
+ 1;
331 n
= splay_tree_lookup (&mm
->splay_tree
, &cur_node
);
334 /* Could be possibly zero size array section. */
336 n
= splay_tree_lookup (&mm
->splay_tree
, &cur_node
);
339 cur_node
.host_start
--;
340 n
= splay_tree_lookup (&mm
->splay_tree
, &cur_node
);
341 cur_node
.host_start
++;
345 gomp_fatal ("Pointer target of array section "
347 cur_node
.host_start
-= n
->host_start
;
348 cur_node
.tgt_offset
= n
->tgt
->tgt_start
+ n
->tgt_offset
349 + cur_node
.host_start
;
350 /* At this point tgt_offset is target address of the
351 array section. Now subtract bias to get what we want
352 to initialize the pointer with. */
353 cur_node
.tgt_offset
-= sizes
[i
];
354 /* FIXME: see above FIXME comment. */
355 devicep
->host2dev_func (devicep
->target_id
,
356 (void *) (tgt
->tgt_start
358 (void *) &cur_node
.tgt_offset
,
361 case GOMP_MAP_TO_PSET
:
362 /* FIXME: see above FIXME comment. */
363 devicep
->host2dev_func (devicep
->target_id
,
364 (void *) (tgt
->tgt_start
366 (void *) k
->host_start
,
367 k
->host_end
- k
->host_start
);
369 for (j
= i
+ 1; j
< mapnum
; j
++)
370 if (!GOMP_MAP_POINTER_P (get_kind (is_openacc
, kinds
, j
)
373 else if ((uintptr_t) hostaddrs
[j
] < k
->host_start
374 || ((uintptr_t) hostaddrs
[j
] + sizeof (void *)
382 = (uintptr_t) *(void **) hostaddrs
[j
];
383 if (cur_node
.host_start
== (uintptr_t) NULL
)
385 cur_node
.tgt_offset
= (uintptr_t) NULL
;
386 /* FIXME: see above FIXME comment. */
387 devicep
->host2dev_func (devicep
->target_id
,
388 (void *) (tgt
->tgt_start
+ k
->tgt_offset
389 + ((uintptr_t) hostaddrs
[j
]
391 (void *) &cur_node
.tgt_offset
,
396 /* Add bias to the pointer value. */
397 cur_node
.host_start
+= sizes
[j
];
398 cur_node
.host_end
= cur_node
.host_start
+ 1;
399 n
= splay_tree_lookup (&mm
->splay_tree
, &cur_node
);
402 /* Could be possibly zero size array section. */
404 n
= splay_tree_lookup (&mm
->splay_tree
,
408 cur_node
.host_start
--;
409 n
= splay_tree_lookup (&mm
->splay_tree
,
411 cur_node
.host_start
++;
415 gomp_fatal ("Pointer target of array section "
417 cur_node
.host_start
-= n
->host_start
;
418 cur_node
.tgt_offset
= n
->tgt
->tgt_start
420 + cur_node
.host_start
;
421 /* At this point tgt_offset is target address of the
422 array section. Now subtract bias to get what we
423 want to initialize the pointer with. */
424 cur_node
.tgt_offset
-= sizes
[j
];
425 /* FIXME: see above FIXME comment. */
426 devicep
->host2dev_func (devicep
->target_id
,
427 (void *) (tgt
->tgt_start
+ k
->tgt_offset
428 + ((uintptr_t) hostaddrs
[j
]
430 (void *) &cur_node
.tgt_offset
,
435 case GOMP_MAP_FORCE_PRESENT
:
437 /* We already looked up the memory region above and it
439 size_t size
= k
->host_end
- k
->host_start
;
440 gomp_fatal ("present clause: !acc_is_present (%p, "
441 "%zd (0x%zx))", (void *) k
->host_start
,
445 case GOMP_MAP_FORCE_DEVICEPTR
:
446 assert (k
->host_end
- k
->host_start
== sizeof (void *));
448 devicep
->host2dev_func (devicep
->target_id
,
449 (void *) (tgt
->tgt_start
451 (void *) k
->host_start
,
455 gomp_fatal ("%s: unhandled kind 0x%.2x", __FUNCTION__
,
465 for (i
= 0; i
< mapnum
; i
++)
467 if (tgt
->list
[i
] == NULL
)
468 cur_node
.tgt_offset
= (uintptr_t) NULL
;
470 cur_node
.tgt_offset
= tgt
->list
[i
]->tgt
->tgt_start
471 + tgt
->list
[i
]->tgt_offset
;
472 /* FIXME: see above FIXME comment. */
473 devicep
->host2dev_func (devicep
->target_id
,
474 (void *) (tgt
->tgt_start
475 + i
* sizeof (void *)),
476 (void *) &cur_node
.tgt_offset
,
481 gomp_mutex_unlock (&mm
->lock
);
486 gomp_unmap_tgt (struct target_mem_desc
*tgt
)
488 /* Deallocate on target the tgt->tgt_start .. tgt->tgt_end region. */
490 tgt
->device_descr
->free_func (tgt
->device_descr
->target_id
, tgt
->to_free
);
496 /* Decrease the refcount for a set of mapped variables, and queue asychronous
497 copies from the device back to the host after any work that has been issued.
498 Because the regions are still "live", increment an asynchronous reference
499 count to indicate that they should not be unmapped from host-side data
500 structures until the asynchronous copy has completed. */
502 attribute_hidden
void
503 gomp_copy_from_async (struct target_mem_desc
*tgt
)
505 struct gomp_device_descr
*devicep
= tgt
->device_descr
;
506 struct gomp_memory_mapping
*mm
= tgt
->mem_map
;
509 gomp_mutex_lock (&mm
->lock
);
511 for (i
= 0; i
< tgt
->list_count
; i
++)
512 if (tgt
->list
[i
] == NULL
)
514 else if (tgt
->list
[i
]->refcount
> 1)
516 tgt
->list
[i
]->refcount
--;
517 tgt
->list
[i
]->async_refcount
++;
521 splay_tree_key k
= tgt
->list
[i
];
523 devicep
->dev2host_func (devicep
->target_id
, (void *) k
->host_start
,
524 (void *) (k
->tgt
->tgt_start
+ k
->tgt_offset
),
525 k
->host_end
- k
->host_start
);
528 gomp_mutex_unlock (&mm
->lock
);
531 /* Unmap variables described by TGT. If DO_COPYFROM is true, copy relevant
532 variables back from device to host: if it is false, it is assumed that this
533 has been done already, i.e. by gomp_copy_from_async above. */
535 attribute_hidden
void
536 gomp_unmap_vars (struct target_mem_desc
*tgt
, bool do_copyfrom
)
538 struct gomp_device_descr
*devicep
= tgt
->device_descr
;
539 struct gomp_memory_mapping
*mm
= tgt
->mem_map
;
541 if (tgt
->list_count
== 0)
547 gomp_mutex_lock (&mm
->lock
);
550 for (i
= 0; i
< tgt
->list_count
; i
++)
551 if (tgt
->list
[i
] == NULL
)
553 else if (tgt
->list
[i
]->refcount
> 1)
554 tgt
->list
[i
]->refcount
--;
555 else if (tgt
->list
[i
]->async_refcount
> 0)
556 tgt
->list
[i
]->async_refcount
--;
559 splay_tree_key k
= tgt
->list
[i
];
560 if (k
->copy_from
&& do_copyfrom
)
561 devicep
->dev2host_func (devicep
->target_id
, (void *) k
->host_start
,
562 (void *) (k
->tgt
->tgt_start
+ k
->tgt_offset
),
563 k
->host_end
- k
->host_start
);
564 splay_tree_remove (&mm
->splay_tree
, k
);
565 if (k
->tgt
->refcount
> 1)
568 gomp_unmap_tgt (k
->tgt
);
571 if (tgt
->refcount
> 1)
574 gomp_unmap_tgt (tgt
);
576 gomp_mutex_unlock (&mm
->lock
);
580 gomp_update (struct gomp_device_descr
*devicep
, struct gomp_memory_mapping
*mm
,
581 size_t mapnum
, void **hostaddrs
, size_t *sizes
, void *kinds
,
585 struct splay_tree_key_s cur_node
;
586 const int typemask
= is_openacc
? 0xff : 0x7;
594 gomp_mutex_lock (&mm
->lock
);
595 for (i
= 0; i
< mapnum
; i
++)
598 cur_node
.host_start
= (uintptr_t) hostaddrs
[i
];
599 cur_node
.host_end
= cur_node
.host_start
+ sizes
[i
];
600 splay_tree_key n
= splay_tree_lookup (&mm
->splay_tree
,
604 int kind
= get_kind (is_openacc
, kinds
, i
);
605 if (n
->host_start
> cur_node
.host_start
606 || n
->host_end
< cur_node
.host_end
)
607 gomp_fatal ("Trying to update [%p..%p) object when"
608 "only [%p..%p) is mapped",
609 (void *) cur_node
.host_start
,
610 (void *) cur_node
.host_end
,
611 (void *) n
->host_start
,
612 (void *) n
->host_end
);
613 if (GOMP_MAP_COPY_TO_P (kind
& typemask
))
614 devicep
->host2dev_func (devicep
->target_id
,
615 (void *) (n
->tgt
->tgt_start
617 + cur_node
.host_start
619 (void *) cur_node
.host_start
,
620 cur_node
.host_end
- cur_node
.host_start
);
621 if (GOMP_MAP_COPY_FROM_P (kind
& typemask
))
622 devicep
->dev2host_func (devicep
->target_id
,
623 (void *) cur_node
.host_start
,
624 (void *) (n
->tgt
->tgt_start
626 + cur_node
.host_start
628 cur_node
.host_end
- cur_node
.host_start
);
631 gomp_fatal ("Trying to update [%p..%p) object that is not mapped",
632 (void *) cur_node
.host_start
,
633 (void *) cur_node
.host_end
);
635 gomp_mutex_unlock (&mm
->lock
);
638 /* This function should be called from every offload image.
639 It gets the descriptor of the host func and var tables HOST_TABLE, TYPE of
640 the target, and TARGET_DATA needed by target plugin. */
643 GOMP_offload_register (void *host_table
, enum offload_target_type target_type
,
646 offload_images
= gomp_realloc (offload_images
,
647 (num_offload_images
+ 1)
648 * sizeof (struct offload_image_descr
));
650 offload_images
[num_offload_images
].type
= target_type
;
651 offload_images
[num_offload_images
].host_table
= host_table
;
652 offload_images
[num_offload_images
].target_data
= target_data
;
654 num_offload_images
++;
657 /* This function initializes the target device, specified by DEVICEP. DEVICEP
658 must be locked on entry, and remains locked on return. */
660 attribute_hidden
void
661 gomp_init_device (struct gomp_device_descr
*devicep
)
663 devicep
->init_device_func (devicep
->target_id
);
664 devicep
->is_initialized
= true;
667 /* Initialize address mapping tables. MM must be locked on entry, and remains
670 attribute_hidden
void
671 gomp_init_tables (struct gomp_device_descr
*devicep
,
672 struct gomp_memory_mapping
*mm
)
674 /* Get address mapping table for device. */
675 struct mapping_table
*table
= NULL
;
676 int num_entries
= devicep
->get_table_func (devicep
->target_id
, &table
);
678 /* Insert host-target address mapping into dev_splay_tree. */
680 for (i
= 0; i
< num_entries
; i
++)
682 struct target_mem_desc
*tgt
= gomp_malloc (sizeof (*tgt
));
684 tgt
->array
= gomp_malloc (sizeof (*tgt
->array
));
685 tgt
->tgt_start
= table
[i
].tgt_start
;
686 tgt
->tgt_end
= table
[i
].tgt_end
;
689 tgt
->device_descr
= devicep
;
690 splay_tree_node node
= tgt
->array
;
691 splay_tree_key k
= &node
->key
;
692 k
->host_start
= table
[i
].host_start
;
693 k
->host_end
= table
[i
].host_end
;
696 k
->copy_from
= false;
700 splay_tree_insert (&mm
->splay_tree
, node
);
704 mm
->is_initialized
= true;
707 /* Free address mapping tables. MM must be locked on entry, and remains locked
710 attribute_hidden
void
711 gomp_free_memmap (struct gomp_memory_mapping
*mm
)
713 while (mm
->splay_tree
.root
)
715 struct target_mem_desc
*tgt
= mm
->splay_tree
.root
->key
.tgt
;
717 splay_tree_remove (&mm
->splay_tree
, &mm
->splay_tree
.root
->key
);
722 mm
->is_initialized
= false;
725 /* This function de-initializes the target device, specified by DEVICEP.
726 DEVICEP must be locked on entry, and remains locked on return. */
728 attribute_hidden
void
729 gomp_fini_device (struct gomp_device_descr
*devicep
)
731 if (devicep
->is_initialized
)
732 devicep
->fini_device_func (devicep
->target_id
);
734 devicep
->is_initialized
= false;
737 /* Called when encountering a target directive. If DEVICE
738 is GOMP_DEVICE_ICV, it means use device-var ICV. If it is
739 GOMP_DEVICE_HOST_FALLBACK (or any value
740 larger than last available hw device), use host fallback.
741 FN is address of host code, OFFLOAD_TABLE contains value of the
742 __OFFLOAD_TABLE__ symbol in the shared library or binary that invokes
743 GOMP_target. HOSTADDRS, SIZES and KINDS are arrays
744 with MAPNUM entries, with addresses of the host objects,
745 sizes of the host objects (resp. for pointer kind pointer bias
746 and assumed sizeof (void *) size) and kinds. */
749 GOMP_target (int device
, void (*fn
) (void *), const void *offload_table
,
750 size_t mapnum
, void **hostaddrs
, size_t *sizes
,
751 unsigned char *kinds
)
753 struct gomp_device_descr
*devicep
= resolve_device (device
);
756 || !(devicep
->capabilities
& GOMP_OFFLOAD_CAP_OPENMP_400
))
759 struct gomp_thread old_thr
, *thr
= gomp_thread ();
761 memset (thr
, '\0', sizeof (*thr
));
762 if (gomp_places_list
)
764 thr
->place
= old_thr
.place
;
765 thr
->ts
.place_partition_len
= gomp_places_list_len
;
768 gomp_free_thread (thr
);
773 gomp_mutex_lock (&devicep
->lock
);
774 if (!devicep
->is_initialized
)
775 gomp_init_device (devicep
);
776 gomp_mutex_unlock (&devicep
->lock
);
780 if (devicep
->capabilities
& GOMP_OFFLOAD_CAP_NATIVE_EXEC
)
781 fn_addr
= (void *) fn
;
784 struct gomp_memory_mapping
*mm
= &devicep
->mem_map
;
785 gomp_mutex_lock (&mm
->lock
);
787 if (!mm
->is_initialized
)
788 gomp_init_tables (devicep
, mm
);
790 struct splay_tree_key_s k
;
791 k
.host_start
= (uintptr_t) fn
;
792 k
.host_end
= k
.host_start
+ 1;
793 splay_tree_key tgt_fn
= splay_tree_lookup (&mm
->splay_tree
, &k
);
795 gomp_fatal ("Target function wasn't mapped");
797 gomp_mutex_unlock (&mm
->lock
);
799 fn_addr
= (void *) tgt_fn
->tgt
->tgt_start
;
802 struct target_mem_desc
*tgt_vars
803 = gomp_map_vars (devicep
, mapnum
, hostaddrs
, NULL
, sizes
, kinds
, false,
805 struct gomp_thread old_thr
, *thr
= gomp_thread ();
807 memset (thr
, '\0', sizeof (*thr
));
808 if (gomp_places_list
)
810 thr
->place
= old_thr
.place
;
811 thr
->ts
.place_partition_len
= gomp_places_list_len
;
813 devicep
->run_func (devicep
->target_id
, fn_addr
, (void *) tgt_vars
->tgt_start
);
814 gomp_free_thread (thr
);
816 gomp_unmap_vars (tgt_vars
, true);
820 GOMP_target_data (int device
, const void *offload_table
, size_t mapnum
,
821 void **hostaddrs
, size_t *sizes
, unsigned char *kinds
)
823 struct gomp_device_descr
*devicep
= resolve_device (device
);
826 || !(devicep
->capabilities
& GOMP_OFFLOAD_CAP_OPENMP_400
))
829 struct gomp_task_icv
*icv
= gomp_icv (false);
830 if (icv
->target_data
)
832 /* Even when doing a host fallback, if there are any active
833 #pragma omp target data constructs, need to remember the
834 new #pragma omp target data, otherwise GOMP_target_end_data
835 would get out of sync. */
836 struct target_mem_desc
*tgt
837 = gomp_map_vars (NULL
, 0, NULL
, NULL
, NULL
, NULL
, false, false);
838 tgt
->prev
= icv
->target_data
;
839 icv
->target_data
= tgt
;
844 gomp_mutex_lock (&devicep
->lock
);
845 if (!devicep
->is_initialized
)
846 gomp_init_device (devicep
);
847 gomp_mutex_unlock (&devicep
->lock
);
849 struct gomp_memory_mapping
*mm
= &devicep
->mem_map
;
850 gomp_mutex_lock (&mm
->lock
);
851 if (!mm
->is_initialized
)
852 gomp_init_tables (devicep
, mm
);
853 gomp_mutex_unlock (&mm
->lock
);
855 struct target_mem_desc
*tgt
856 = gomp_map_vars (devicep
, mapnum
, hostaddrs
, NULL
, sizes
, kinds
, false,
858 struct gomp_task_icv
*icv
= gomp_icv (true);
859 tgt
->prev
= icv
->target_data
;
860 icv
->target_data
= tgt
;
864 GOMP_target_end_data (void)
866 struct gomp_task_icv
*icv
= gomp_icv (false);
867 if (icv
->target_data
)
869 struct target_mem_desc
*tgt
= icv
->target_data
;
870 icv
->target_data
= tgt
->prev
;
871 gomp_unmap_vars (tgt
, true);
876 GOMP_target_update (int device
, const void *offload_table
, size_t mapnum
,
877 void **hostaddrs
, size_t *sizes
, unsigned char *kinds
)
879 struct gomp_device_descr
*devicep
= resolve_device (device
);
882 || !(devicep
->capabilities
& GOMP_OFFLOAD_CAP_OPENMP_400
))
885 gomp_mutex_lock (&devicep
->lock
);
886 if (!devicep
->is_initialized
)
887 gomp_init_device (devicep
);
888 gomp_mutex_unlock (&devicep
->lock
);
890 struct gomp_memory_mapping
*mm
= &devicep
->mem_map
;
891 gomp_mutex_lock (&mm
->lock
);
892 if (!mm
->is_initialized
)
893 gomp_init_tables (devicep
, mm
);
894 gomp_mutex_unlock (&mm
->lock
);
896 gomp_update (devicep
, mm
, mapnum
, hostaddrs
, sizes
, kinds
, false);
900 GOMP_teams (unsigned int num_teams
, unsigned int thread_limit
)
904 struct gomp_task_icv
*icv
= gomp_icv (true);
905 icv
->thread_limit_var
906 = thread_limit
> INT_MAX
? UINT_MAX
: thread_limit
;
911 #ifdef PLUGIN_SUPPORT
913 /* This function tries to load a plugin for DEVICE. Name of plugin is passed
915 The handles of the found functions are stored in the corresponding fields
916 of DEVICE. The function returns TRUE on success and FALSE otherwise. */
919 gomp_load_plugin_for_device (struct gomp_device_descr
*device
,
920 const char *plugin_name
)
922 char *err
= NULL
, *last_missing
= NULL
;
923 int optional_present
, optional_total
;
925 /* Clear any existing error. */
928 void *plugin_handle
= dlopen (plugin_name
, RTLD_LAZY
);
935 /* Check if all required functions are available in the plugin and store
940 device->f##_func = dlsym (plugin_handle, "GOMP_OFFLOAD_" #f); \
946 /* Similar, but missing functions are not an error. */
947 #define DLSYM_OPT(f, n) \
951 device->f##_func = dlsym (plugin_handle, "GOMP_OFFLOAD_" #n); \
952 tmp_err = dlerror (); \
953 if (tmp_err == NULL) \
954 optional_present++; \
964 DLSYM (get_num_devices
);
965 DLSYM (register_image
);
973 device
->capabilities
= device
->get_caps_func ();
974 if (device
->capabilities
& GOMP_OFFLOAD_CAP_OPENMP_400
)
976 if (device
->capabilities
& GOMP_OFFLOAD_CAP_OPENACC_200
)
978 optional_present
= optional_total
= 0;
979 DLSYM_OPT (openacc
.exec
, openacc_parallel
);
980 DLSYM_OPT (openacc
.open_device
, openacc_open_device
);
981 DLSYM_OPT (openacc
.close_device
, openacc_close_device
);
982 DLSYM_OPT (openacc
.get_device_num
, openacc_get_device_num
);
983 DLSYM_OPT (openacc
.set_device_num
, openacc_set_device_num
);
984 DLSYM_OPT (openacc
.register_async_cleanup
,
985 openacc_register_async_cleanup
);
986 DLSYM_OPT (openacc
.async_test
, openacc_async_test
);
987 DLSYM_OPT (openacc
.async_test_all
, openacc_async_test_all
);
988 DLSYM_OPT (openacc
.async_wait
, openacc_async_wait
);
989 DLSYM_OPT (openacc
.async_wait_async
, openacc_async_wait_async
);
990 DLSYM_OPT (openacc
.async_wait_all
, openacc_async_wait_all
);
991 DLSYM_OPT (openacc
.async_wait_all_async
, openacc_async_wait_all_async
);
992 DLSYM_OPT (openacc
.async_set_async
, openacc_async_set_async
);
993 DLSYM_OPT (openacc
.create_thread_data
, openacc_create_thread_data
);
994 DLSYM_OPT (openacc
.destroy_thread_data
, openacc_destroy_thread_data
);
995 /* Require all the OpenACC handlers if we have
996 GOMP_OFFLOAD_CAP_OPENACC_200. */
997 if (optional_present
!= optional_total
)
999 err
= "plugin missing OpenACC handler function";
1002 optional_present
= optional_total
= 0;
1003 DLSYM_OPT (openacc
.cuda
.get_current_device
,
1004 openacc_get_current_cuda_device
);
1005 DLSYM_OPT (openacc
.cuda
.get_current_context
,
1006 openacc_get_current_cuda_context
);
1007 DLSYM_OPT (openacc
.cuda
.get_stream
, openacc_get_cuda_stream
);
1008 DLSYM_OPT (openacc
.cuda
.set_stream
, openacc_set_cuda_stream
);
1009 /* Make sure all the CUDA functions are there if any of them are. */
1010 if (optional_present
&& optional_present
!= optional_total
)
1012 err
= "plugin missing OpenACC CUDA handler function";
1022 gomp_error ("while loading %s: %s", plugin_name
, err
);
1024 gomp_error ("missing function was %s", last_missing
);
1026 dlclose (plugin_handle
);
1031 /* This function adds a compatible offload image IMAGE to an accelerator device
1032 DEVICE. DEVICE must be locked on entry, and remains locked on return. */
1035 gomp_register_image_for_device (struct gomp_device_descr
*device
,
1036 struct offload_image_descr
*image
)
1038 if (!device
->offload_regions_registered
1039 && (device
->type
== image
->type
1040 || device
->type
== OFFLOAD_TARGET_TYPE_HOST
))
1042 device
->register_image_func (image
->host_table
, image
->target_data
);
1043 device
->offload_regions_registered
= true;
1047 /* This function initializes the runtime needed for offloading.
1048 It parses the list of offload targets and tries to load the plugins for
1049 these targets. On return, the variables NUM_DEVICES and NUM_DEVICES_OPENMP
1050 will be set, and the array DEVICES initialized, containing descriptors for
1051 corresponding devices, first the GOMP_OFFLOAD_CAP_OPENMP_400 ones, follows
1055 gomp_target_init (void)
1057 const char *prefix
="libgomp-plugin-";
1058 const char *suffix
= ".so.1";
1059 const char *cur
, *next
;
1061 int i
, new_num_devices
;
1066 cur
= OFFLOAD_TARGETS
;
1070 struct gomp_device_descr current_device
;
1072 next
= strchr (cur
, ',');
1074 plugin_name
= (char *) malloc (1 + (next
? next
- cur
: strlen (cur
))
1075 + strlen (prefix
) + strlen (suffix
));
1082 strcpy (plugin_name
, prefix
);
1083 strncat (plugin_name
, cur
, next
? next
- cur
: strlen (cur
));
1084 strcat (plugin_name
, suffix
);
1086 if (gomp_load_plugin_for_device (¤t_device
, plugin_name
))
1088 new_num_devices
= current_device
.get_num_devices_func ();
1089 if (new_num_devices
>= 1)
1091 /* Augment DEVICES and NUM_DEVICES. */
1093 devices
= realloc (devices
, (num_devices
+ new_num_devices
)
1094 * sizeof (struct gomp_device_descr
));
1102 current_device
.name
= current_device
.get_name_func ();
1103 /* current_device.capabilities has already been set. */
1104 current_device
.type
= current_device
.get_type_func ();
1105 current_device
.mem_map
.is_initialized
= false;
1106 current_device
.mem_map
.splay_tree
.root
= NULL
;
1107 current_device
.is_initialized
= false;
1108 current_device
.offload_regions_registered
= false;
1109 current_device
.openacc
.data_environ
= NULL
;
1110 current_device
.openacc
.target_data
= NULL
;
1111 for (i
= 0; i
< new_num_devices
; i
++)
1113 current_device
.target_id
= i
;
1114 devices
[num_devices
] = current_device
;
1115 gomp_mutex_init (&devices
[num_devices
].mem_map
.lock
);
1116 gomp_mutex_init (&devices
[num_devices
].lock
);
1127 /* In DEVICES, sort the GOMP_OFFLOAD_CAP_OPENMP_400 ones first, and set
1128 NUM_DEVICES_OPENMP. */
1129 struct gomp_device_descr
*devices_s
1130 = malloc (num_devices
* sizeof (struct gomp_device_descr
));
1137 num_devices_openmp
= 0;
1138 for (i
= 0; i
< num_devices
; i
++)
1139 if (devices
[i
].capabilities
& GOMP_OFFLOAD_CAP_OPENMP_400
)
1140 devices_s
[num_devices_openmp
++] = devices
[i
];
1141 int num_devices_after_openmp
= num_devices_openmp
;
1142 for (i
= 0; i
< num_devices
; i
++)
1143 if (!(devices
[i
].capabilities
& GOMP_OFFLOAD_CAP_OPENMP_400
))
1144 devices_s
[num_devices_after_openmp
++] = devices
[i
];
1146 devices
= devices_s
;
1148 for (i
= 0; i
< num_devices
; i
++)
1152 for (j
= 0; j
< num_offload_images
; j
++)
1153 gomp_register_image_for_device (&devices
[i
], &offload_images
[j
]);
1155 /* The 'devices' array can be moved (by the realloc call) until we have
1156 found all the plugins, so registering with the OpenACC runtime (which
1157 takes a copy of the pointer argument) must be delayed until now. */
1158 if (devices
[i
].capabilities
& GOMP_OFFLOAD_CAP_OPENACC_200
)
1159 goacc_register (&devices
[i
]);
1162 free (offload_images
);
1163 offload_images
= NULL
;
1164 num_offload_images
= 0;
1167 #else /* PLUGIN_SUPPORT */
1168 /* If dlfcn.h is unavailable we always fallback to host execution.
1169 GOMP_target* routines are just stubs for this case. */
1171 gomp_target_init (void)
1174 #endif /* PLUGIN_SUPPORT */