1 /* Copyright (C) 2013-2016 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"
36 #ifdef HAVE_INTTYPES_H
37 # include <inttypes.h> /* For PRIu64. */
45 #include "plugin-suffix.h"
48 static void gomp_target_init (void);
50 /* The whole initialization code for offloading plugins is only run one. */
51 static pthread_once_t gomp_is_initialized
= PTHREAD_ONCE_INIT
;
53 /* Mutex for offload image registration. */
54 static gomp_mutex_t register_lock
;
56 /* This structure describes an offload image.
57 It contains type of the target device, pointer to host table descriptor, and
58 pointer to target data. */
59 struct offload_image_descr
{
61 enum offload_target_type type
;
62 const void *host_table
;
63 const void *target_data
;
66 /* Array of descriptors of offload images. */
67 static struct offload_image_descr
*offload_images
;
69 /* Total number of offload images. */
70 static int num_offload_images
;
72 /* Array of descriptors for all available devices. */
73 static struct gomp_device_descr
*devices
;
75 /* Total number of available devices. */
76 static int num_devices
;
78 /* Number of GOMP_OFFLOAD_CAP_OPENMP_400 devices. */
79 static int num_devices_openmp
;
81 /* Similar to gomp_realloc, but release register_lock before gomp_fatal. */
84 gomp_realloc_unlock (void *old
, size_t size
)
86 void *ret
= realloc (old
, size
);
89 gomp_mutex_unlock (®ister_lock
);
90 gomp_fatal ("Out of memory allocating %lu bytes", (unsigned long) size
);
96 gomp_init_targets_once (void)
98 (void) pthread_once (&gomp_is_initialized
, gomp_target_init
);
102 gomp_get_num_devices (void)
104 gomp_init_targets_once ();
105 return num_devices_openmp
;
108 static struct gomp_device_descr
*
109 resolve_device (int device_id
)
111 if (device_id
== GOMP_DEVICE_ICV
)
113 struct gomp_task_icv
*icv
= gomp_icv (false);
114 device_id
= icv
->default_device_var
;
117 if (device_id
< 0 || device_id
>= gomp_get_num_devices ())
120 gomp_mutex_lock (&devices
[device_id
].lock
);
121 if (devices
[device_id
].state
== GOMP_DEVICE_UNINITIALIZED
)
122 gomp_init_device (&devices
[device_id
]);
123 else if (devices
[device_id
].state
== GOMP_DEVICE_FINALIZED
)
125 gomp_mutex_unlock (&devices
[device_id
].lock
);
128 gomp_mutex_unlock (&devices
[device_id
].lock
);
130 return &devices
[device_id
];
134 static inline splay_tree_key
135 gomp_map_lookup (splay_tree mem_map
, splay_tree_key key
)
137 if (key
->host_start
!= key
->host_end
)
138 return splay_tree_lookup (mem_map
, key
);
141 splay_tree_key n
= splay_tree_lookup (mem_map
, key
);
146 n
= splay_tree_lookup (mem_map
, key
);
150 return splay_tree_lookup (mem_map
, key
);
153 static inline splay_tree_key
154 gomp_map_0len_lookup (splay_tree mem_map
, splay_tree_key key
)
156 if (key
->host_start
!= key
->host_end
)
157 return splay_tree_lookup (mem_map
, key
);
160 splay_tree_key n
= splay_tree_lookup (mem_map
, key
);
165 /* Handle the case where gomp_map_lookup, splay_tree_lookup or
166 gomp_map_0len_lookup found oldn for newn.
167 Helper function of gomp_map_vars. */
170 gomp_map_vars_existing (struct gomp_device_descr
*devicep
, splay_tree_key oldn
,
171 splay_tree_key newn
, struct target_var_desc
*tgt_var
,
175 tgt_var
->copy_from
= GOMP_MAP_COPY_FROM_P (kind
);
176 tgt_var
->always_copy_from
= GOMP_MAP_ALWAYS_FROM_P (kind
);
177 tgt_var
->offset
= newn
->host_start
- oldn
->host_start
;
178 tgt_var
->length
= newn
->host_end
- newn
->host_start
;
180 if ((kind
& GOMP_MAP_FLAG_FORCE
)
181 || oldn
->host_start
> newn
->host_start
182 || oldn
->host_end
< newn
->host_end
)
184 gomp_mutex_unlock (&devicep
->lock
);
185 gomp_fatal ("Trying to map into device [%p..%p) object when "
186 "[%p..%p) is already mapped",
187 (void *) newn
->host_start
, (void *) newn
->host_end
,
188 (void *) oldn
->host_start
, (void *) oldn
->host_end
);
191 if (GOMP_MAP_ALWAYS_TO_P (kind
))
192 devicep
->host2dev_func (devicep
->target_id
,
193 (void *) (oldn
->tgt
->tgt_start
+ oldn
->tgt_offset
194 + newn
->host_start
- oldn
->host_start
),
195 (void *) newn
->host_start
,
196 newn
->host_end
- newn
->host_start
);
197 if (oldn
->refcount
!= REFCOUNT_INFINITY
)
202 get_kind (bool short_mapkind
, void *kinds
, int idx
)
204 return short_mapkind
? ((unsigned short *) kinds
)[idx
]
205 : ((unsigned char *) kinds
)[idx
];
209 gomp_map_pointer (struct target_mem_desc
*tgt
, uintptr_t host_ptr
,
210 uintptr_t target_offset
, uintptr_t bias
)
212 struct gomp_device_descr
*devicep
= tgt
->device_descr
;
213 struct splay_tree_s
*mem_map
= &devicep
->mem_map
;
214 struct splay_tree_key_s cur_node
;
216 cur_node
.host_start
= host_ptr
;
217 if (cur_node
.host_start
== (uintptr_t) NULL
)
219 cur_node
.tgt_offset
= (uintptr_t) NULL
;
220 /* FIXME: see comment about coalescing host/dev transfers below. */
221 devicep
->host2dev_func (devicep
->target_id
,
222 (void *) (tgt
->tgt_start
+ target_offset
),
223 (void *) &cur_node
.tgt_offset
,
227 /* Add bias to the pointer value. */
228 cur_node
.host_start
+= bias
;
229 cur_node
.host_end
= cur_node
.host_start
;
230 splay_tree_key n
= gomp_map_lookup (mem_map
, &cur_node
);
233 gomp_mutex_unlock (&devicep
->lock
);
234 gomp_fatal ("Pointer target of array section wasn't mapped");
236 cur_node
.host_start
-= n
->host_start
;
238 = n
->tgt
->tgt_start
+ n
->tgt_offset
+ cur_node
.host_start
;
239 /* At this point tgt_offset is target address of the
240 array section. Now subtract bias to get what we want
241 to initialize the pointer with. */
242 cur_node
.tgt_offset
-= bias
;
243 /* FIXME: see comment about coalescing host/dev transfers below. */
244 devicep
->host2dev_func (devicep
->target_id
,
245 (void *) (tgt
->tgt_start
+ target_offset
),
246 (void *) &cur_node
.tgt_offset
,
251 gomp_map_fields_existing (struct target_mem_desc
*tgt
, splay_tree_key n
,
252 size_t first
, size_t i
, void **hostaddrs
,
253 size_t *sizes
, void *kinds
)
255 struct gomp_device_descr
*devicep
= tgt
->device_descr
;
256 struct splay_tree_s
*mem_map
= &devicep
->mem_map
;
257 struct splay_tree_key_s cur_node
;
259 const bool short_mapkind
= true;
260 const int typemask
= short_mapkind
? 0xff : 0x7;
262 cur_node
.host_start
= (uintptr_t) hostaddrs
[i
];
263 cur_node
.host_end
= cur_node
.host_start
+ sizes
[i
];
264 splay_tree_key n2
= splay_tree_lookup (mem_map
, &cur_node
);
265 kind
= get_kind (short_mapkind
, kinds
, i
);
268 && n2
->host_start
- n
->host_start
== n2
->tgt_offset
- n
->tgt_offset
)
270 gomp_map_vars_existing (devicep
, n2
, &cur_node
,
271 &tgt
->list
[i
], kind
& typemask
);
276 if (cur_node
.host_start
> (uintptr_t) hostaddrs
[first
- 1])
278 cur_node
.host_start
--;
279 n2
= splay_tree_lookup (mem_map
, &cur_node
);
280 cur_node
.host_start
++;
283 && n2
->host_start
- n
->host_start
284 == n2
->tgt_offset
- n
->tgt_offset
)
286 gomp_map_vars_existing (devicep
, n2
, &cur_node
, &tgt
->list
[i
],
292 n2
= splay_tree_lookup (mem_map
, &cur_node
);
296 && n2
->host_start
- n
->host_start
== n2
->tgt_offset
- n
->tgt_offset
)
298 gomp_map_vars_existing (devicep
, n2
, &cur_node
, &tgt
->list
[i
],
303 gomp_mutex_unlock (&devicep
->lock
);
304 gomp_fatal ("Trying to map into device [%p..%p) structure element when "
305 "other mapped elements from the same structure weren't mapped "
306 "together with it", (void *) cur_node
.host_start
,
307 (void *) cur_node
.host_end
);
310 static inline uintptr_t
311 gomp_map_val (struct target_mem_desc
*tgt
, void **hostaddrs
, size_t i
)
313 if (tgt
->list
[i
].key
!= NULL
)
314 return tgt
->list
[i
].key
->tgt
->tgt_start
315 + tgt
->list
[i
].key
->tgt_offset
316 + tgt
->list
[i
].offset
;
317 if (tgt
->list
[i
].offset
== ~(uintptr_t) 0)
318 return (uintptr_t) hostaddrs
[i
];
319 if (tgt
->list
[i
].offset
== ~(uintptr_t) 1)
321 if (tgt
->list
[i
].offset
== ~(uintptr_t) 2)
322 return tgt
->list
[i
+ 1].key
->tgt
->tgt_start
323 + tgt
->list
[i
+ 1].key
->tgt_offset
324 + tgt
->list
[i
+ 1].offset
325 + (uintptr_t) hostaddrs
[i
]
326 - (uintptr_t) hostaddrs
[i
+ 1];
327 return tgt
->tgt_start
+ tgt
->list
[i
].offset
;
330 attribute_hidden
struct target_mem_desc
*
331 gomp_map_vars (struct gomp_device_descr
*devicep
, size_t mapnum
,
332 void **hostaddrs
, void **devaddrs
, size_t *sizes
, void *kinds
,
333 bool short_mapkind
, enum gomp_map_vars_kind pragma_kind
)
335 size_t i
, tgt_align
, tgt_size
, not_found_cnt
= 0;
336 bool has_firstprivate
= false;
337 const int rshift
= short_mapkind
? 8 : 3;
338 const int typemask
= short_mapkind
? 0xff : 0x7;
339 struct splay_tree_s
*mem_map
= &devicep
->mem_map
;
340 struct splay_tree_key_s cur_node
;
341 struct target_mem_desc
*tgt
342 = gomp_malloc (sizeof (*tgt
) + sizeof (tgt
->list
[0]) * mapnum
);
343 tgt
->list_count
= mapnum
;
344 tgt
->refcount
= pragma_kind
== GOMP_MAP_VARS_ENTER_DATA
? 0 : 1;
345 tgt
->device_descr
= devicep
;
354 tgt_align
= sizeof (void *);
356 if (pragma_kind
== GOMP_MAP_VARS_TARGET
)
358 size_t align
= 4 * sizeof (void *);
360 tgt_size
= mapnum
* sizeof (void *);
363 gomp_mutex_lock (&devicep
->lock
);
364 if (devicep
->state
== GOMP_DEVICE_FINALIZED
)
366 gomp_mutex_unlock (&devicep
->lock
);
371 for (i
= 0; i
< mapnum
; i
++)
373 int kind
= get_kind (short_mapkind
, kinds
, i
);
374 if (hostaddrs
[i
] == NULL
375 || (kind
& typemask
) == GOMP_MAP_FIRSTPRIVATE_INT
)
377 tgt
->list
[i
].key
= NULL
;
378 tgt
->list
[i
].offset
= ~(uintptr_t) 0;
381 else if ((kind
& typemask
) == GOMP_MAP_USE_DEVICE_PTR
)
383 cur_node
.host_start
= (uintptr_t) hostaddrs
[i
];
384 cur_node
.host_end
= cur_node
.host_start
;
385 splay_tree_key n
= gomp_map_lookup (mem_map
, &cur_node
);
388 gomp_mutex_unlock (&devicep
->lock
);
389 gomp_fatal ("use_device_ptr pointer wasn't mapped");
391 cur_node
.host_start
-= n
->host_start
;
393 = (void *) (n
->tgt
->tgt_start
+ n
->tgt_offset
394 + cur_node
.host_start
);
395 tgt
->list
[i
].key
= NULL
;
396 tgt
->list
[i
].offset
= ~(uintptr_t) 0;
399 else if ((kind
& typemask
) == GOMP_MAP_STRUCT
)
401 size_t first
= i
+ 1;
402 size_t last
= i
+ sizes
[i
];
403 cur_node
.host_start
= (uintptr_t) hostaddrs
[i
];
404 cur_node
.host_end
= (uintptr_t) hostaddrs
[last
]
406 tgt
->list
[i
].key
= NULL
;
407 tgt
->list
[i
].offset
= ~(uintptr_t) 2;
408 splay_tree_key n
= splay_tree_lookup (mem_map
, &cur_node
);
411 size_t align
= (size_t) 1 << (kind
>> rshift
);
412 if (tgt_align
< align
)
414 tgt_size
-= (uintptr_t) hostaddrs
[first
]
415 - (uintptr_t) hostaddrs
[i
];
416 tgt_size
= (tgt_size
+ align
- 1) & ~(align
- 1);
417 tgt_size
+= cur_node
.host_end
- (uintptr_t) hostaddrs
[i
];
418 not_found_cnt
+= last
- i
;
419 for (i
= first
; i
<= last
; i
++)
420 tgt
->list
[i
].key
= NULL
;
424 for (i
= first
; i
<= last
; i
++)
425 gomp_map_fields_existing (tgt
, n
, first
, i
, hostaddrs
,
430 else if ((kind
& typemask
) == GOMP_MAP_ALWAYS_POINTER
)
432 tgt
->list
[i
].key
= NULL
;
433 tgt
->list
[i
].offset
= ~(uintptr_t) 1;
434 has_firstprivate
= true;
437 cur_node
.host_start
= (uintptr_t) hostaddrs
[i
];
438 if (!GOMP_MAP_POINTER_P (kind
& typemask
))
439 cur_node
.host_end
= cur_node
.host_start
+ sizes
[i
];
441 cur_node
.host_end
= cur_node
.host_start
+ sizeof (void *);
442 if ((kind
& typemask
) == GOMP_MAP_FIRSTPRIVATE
)
444 tgt
->list
[i
].key
= NULL
;
446 size_t align
= (size_t) 1 << (kind
>> rshift
);
447 if (tgt_align
< align
)
449 tgt_size
= (tgt_size
+ align
- 1) & ~(align
- 1);
450 tgt_size
+= cur_node
.host_end
- cur_node
.host_start
;
451 has_firstprivate
= true;
455 if ((kind
& typemask
) == GOMP_MAP_ZERO_LEN_ARRAY_SECTION
)
457 n
= gomp_map_0len_lookup (mem_map
, &cur_node
);
460 tgt
->list
[i
].key
= NULL
;
461 tgt
->list
[i
].offset
= ~(uintptr_t) 1;
466 n
= splay_tree_lookup (mem_map
, &cur_node
);
467 if (n
&& n
->refcount
!= REFCOUNT_LINK
)
468 gomp_map_vars_existing (devicep
, n
, &cur_node
, &tgt
->list
[i
],
472 tgt
->list
[i
].key
= NULL
;
474 size_t align
= (size_t) 1 << (kind
>> rshift
);
476 if (tgt_align
< align
)
478 tgt_size
= (tgt_size
+ align
- 1) & ~(align
- 1);
479 tgt_size
+= cur_node
.host_end
- cur_node
.host_start
;
480 if ((kind
& typemask
) == GOMP_MAP_TO_PSET
)
483 for (j
= i
+ 1; j
< mapnum
; j
++)
484 if (!GOMP_MAP_POINTER_P (get_kind (short_mapkind
, kinds
, j
)
487 else if ((uintptr_t) hostaddrs
[j
] < cur_node
.host_start
488 || ((uintptr_t) hostaddrs
[j
] + sizeof (void *)
489 > cur_node
.host_end
))
493 tgt
->list
[j
].key
= NULL
;
504 gomp_mutex_unlock (&devicep
->lock
);
505 gomp_fatal ("unexpected aggregation");
507 tgt
->to_free
= devaddrs
[0];
508 tgt
->tgt_start
= (uintptr_t) tgt
->to_free
;
509 tgt
->tgt_end
= tgt
->tgt_start
+ sizes
[0];
511 else if (not_found_cnt
|| pragma_kind
== GOMP_MAP_VARS_TARGET
)
513 /* Allocate tgt_align aligned tgt_size block of memory. */
514 /* FIXME: Perhaps change interface to allocate properly aligned
516 tgt
->to_free
= devicep
->alloc_func (devicep
->target_id
,
517 tgt_size
+ tgt_align
- 1);
518 tgt
->tgt_start
= (uintptr_t) tgt
->to_free
;
519 tgt
->tgt_start
= (tgt
->tgt_start
+ tgt_align
- 1) & ~(tgt_align
- 1);
520 tgt
->tgt_end
= tgt
->tgt_start
+ tgt_size
;
530 if (pragma_kind
== GOMP_MAP_VARS_TARGET
)
531 tgt_size
= mapnum
* sizeof (void *);
534 if (not_found_cnt
|| has_firstprivate
)
537 tgt
->array
= gomp_malloc (not_found_cnt
* sizeof (*tgt
->array
));
538 splay_tree_node array
= tgt
->array
;
539 size_t j
, field_tgt_offset
= 0, field_tgt_clear
= ~(size_t) 0;
540 uintptr_t field_tgt_base
= 0;
542 for (i
= 0; i
< mapnum
; i
++)
543 if (tgt
->list
[i
].key
== NULL
)
545 int kind
= get_kind (short_mapkind
, kinds
, i
);
546 if (hostaddrs
[i
] == NULL
)
548 switch (kind
& typemask
)
550 size_t align
, len
, first
, last
;
552 case GOMP_MAP_FIRSTPRIVATE
:
553 align
= (size_t) 1 << (kind
>> rshift
);
554 tgt_size
= (tgt_size
+ align
- 1) & ~(align
- 1);
555 tgt
->list
[i
].offset
= tgt_size
;
557 devicep
->host2dev_func (devicep
->target_id
,
558 (void *) (tgt
->tgt_start
+ tgt_size
),
559 (void *) hostaddrs
[i
], len
);
562 case GOMP_MAP_FIRSTPRIVATE_INT
:
563 case GOMP_MAP_USE_DEVICE_PTR
:
564 case GOMP_MAP_ZERO_LEN_ARRAY_SECTION
:
566 case GOMP_MAP_STRUCT
:
569 cur_node
.host_start
= (uintptr_t) hostaddrs
[i
];
570 cur_node
.host_end
= (uintptr_t) hostaddrs
[last
]
572 if (tgt
->list
[first
].key
!= NULL
)
574 n
= splay_tree_lookup (mem_map
, &cur_node
);
577 size_t align
= (size_t) 1 << (kind
>> rshift
);
578 tgt_size
-= (uintptr_t) hostaddrs
[first
]
579 - (uintptr_t) hostaddrs
[i
];
580 tgt_size
= (tgt_size
+ align
- 1) & ~(align
- 1);
581 tgt_size
+= (uintptr_t) hostaddrs
[first
]
582 - (uintptr_t) hostaddrs
[i
];
583 field_tgt_base
= (uintptr_t) hostaddrs
[first
];
584 field_tgt_offset
= tgt_size
;
585 field_tgt_clear
= last
;
586 tgt_size
+= cur_node
.host_end
587 - (uintptr_t) hostaddrs
[first
];
590 for (i
= first
; i
<= last
; i
++)
591 gomp_map_fields_existing (tgt
, n
, first
, i
, hostaddrs
,
595 case GOMP_MAP_ALWAYS_POINTER
:
596 cur_node
.host_start
= (uintptr_t) hostaddrs
[i
];
597 cur_node
.host_end
= cur_node
.host_start
+ sizeof (void *);
598 n
= splay_tree_lookup (mem_map
, &cur_node
);
600 || n
->host_start
> cur_node
.host_start
601 || n
->host_end
< cur_node
.host_end
)
603 gomp_mutex_unlock (&devicep
->lock
);
604 gomp_fatal ("always pointer not mapped");
606 if ((get_kind (short_mapkind
, kinds
, i
- 1) & typemask
)
607 != GOMP_MAP_ALWAYS_POINTER
)
608 cur_node
.tgt_offset
= gomp_map_val (tgt
, hostaddrs
, i
- 1);
609 if (cur_node
.tgt_offset
)
610 cur_node
.tgt_offset
-= sizes
[i
];
611 devicep
->host2dev_func (devicep
->target_id
,
612 (void *) (n
->tgt
->tgt_start
614 + cur_node
.host_start
616 (void *) &cur_node
.tgt_offset
,
618 cur_node
.tgt_offset
= n
->tgt
->tgt_start
+ n
->tgt_offset
619 + cur_node
.host_start
- n
->host_start
;
624 splay_tree_key k
= &array
->key
;
625 k
->host_start
= (uintptr_t) hostaddrs
[i
];
626 if (!GOMP_MAP_POINTER_P (kind
& typemask
))
627 k
->host_end
= k
->host_start
+ sizes
[i
];
629 k
->host_end
= k
->host_start
+ sizeof (void *);
630 splay_tree_key n
= splay_tree_lookup (mem_map
, k
);
631 if (n
&& n
->refcount
!= REFCOUNT_LINK
)
632 gomp_map_vars_existing (devicep
, n
, k
, &tgt
->list
[i
],
637 if (n
&& n
->refcount
== REFCOUNT_LINK
)
639 /* Replace target address of the pointer with target address
640 of mapped object in the splay tree. */
641 splay_tree_remove (mem_map
, n
);
644 size_t align
= (size_t) 1 << (kind
>> rshift
);
645 tgt
->list
[i
].key
= k
;
647 if (field_tgt_clear
!= ~(size_t) 0)
649 k
->tgt_offset
= k
->host_start
- field_tgt_base
651 if (i
== field_tgt_clear
)
652 field_tgt_clear
= ~(size_t) 0;
656 tgt_size
= (tgt_size
+ align
- 1) & ~(align
- 1);
657 k
->tgt_offset
= tgt_size
;
658 tgt_size
+= k
->host_end
- k
->host_start
;
660 tgt
->list
[i
].copy_from
= GOMP_MAP_COPY_FROM_P (kind
& typemask
);
661 tgt
->list
[i
].always_copy_from
662 = GOMP_MAP_ALWAYS_FROM_P (kind
& typemask
);
663 tgt
->list
[i
].offset
= 0;
664 tgt
->list
[i
].length
= k
->host_end
- k
->host_start
;
666 k
->async_refcount
= 0;
670 splay_tree_insert (mem_map
, array
);
671 switch (kind
& typemask
)
675 case GOMP_MAP_FORCE_ALLOC
:
676 case GOMP_MAP_FORCE_FROM
:
677 case GOMP_MAP_ALWAYS_FROM
:
680 case GOMP_MAP_TOFROM
:
681 case GOMP_MAP_FORCE_TO
:
682 case GOMP_MAP_FORCE_TOFROM
:
683 case GOMP_MAP_ALWAYS_TO
:
684 case GOMP_MAP_ALWAYS_TOFROM
:
685 /* FIXME: Perhaps add some smarts, like if copying
686 several adjacent fields from host to target, use some
687 host buffer to avoid sending each var individually. */
688 devicep
->host2dev_func (devicep
->target_id
,
689 (void *) (tgt
->tgt_start
691 (void *) k
->host_start
,
692 k
->host_end
- k
->host_start
);
694 case GOMP_MAP_POINTER
:
695 gomp_map_pointer (tgt
, (uintptr_t) *(void **) k
->host_start
,
696 k
->tgt_offset
, sizes
[i
]);
698 case GOMP_MAP_TO_PSET
:
699 /* FIXME: see above FIXME comment. */
700 devicep
->host2dev_func (devicep
->target_id
,
701 (void *) (tgt
->tgt_start
703 (void *) k
->host_start
,
704 k
->host_end
- k
->host_start
);
706 for (j
= i
+ 1; j
< mapnum
; j
++)
707 if (!GOMP_MAP_POINTER_P (get_kind (short_mapkind
, kinds
,
711 else if ((uintptr_t) hostaddrs
[j
] < k
->host_start
712 || ((uintptr_t) hostaddrs
[j
] + sizeof (void *)
717 tgt
->list
[j
].key
= k
;
718 tgt
->list
[j
].copy_from
= false;
719 tgt
->list
[j
].always_copy_from
= false;
720 if (k
->refcount
!= REFCOUNT_INFINITY
)
722 gomp_map_pointer (tgt
,
723 (uintptr_t) *(void **) hostaddrs
[j
],
725 + ((uintptr_t) hostaddrs
[j
]
731 case GOMP_MAP_FORCE_PRESENT
:
733 /* We already looked up the memory region above and it
735 size_t size
= k
->host_end
- k
->host_start
;
736 gomp_mutex_unlock (&devicep
->lock
);
737 #ifdef HAVE_INTTYPES_H
738 gomp_fatal ("present clause: !acc_is_present (%p, "
739 "%"PRIu64
" (0x%"PRIx64
"))",
740 (void *) k
->host_start
,
741 (uint64_t) size
, (uint64_t) size
);
743 gomp_fatal ("present clause: !acc_is_present (%p, "
744 "%lu (0x%lx))", (void *) k
->host_start
,
745 (unsigned long) size
, (unsigned long) size
);
749 case GOMP_MAP_FORCE_DEVICEPTR
:
750 assert (k
->host_end
- k
->host_start
== sizeof (void *));
752 devicep
->host2dev_func (devicep
->target_id
,
753 (void *) (tgt
->tgt_start
755 (void *) k
->host_start
,
759 gomp_mutex_unlock (&devicep
->lock
);
760 gomp_fatal ("%s: unhandled kind 0x%.2x", __FUNCTION__
,
766 /* Set link pointer on target to the device address of the
768 void *tgt_addr
= (void *) (tgt
->tgt_start
+ k
->tgt_offset
);
769 devicep
->host2dev_func (devicep
->target_id
,
770 (void *) n
->tgt_offset
,
771 &tgt_addr
, sizeof (void *));
778 if (pragma_kind
== GOMP_MAP_VARS_TARGET
)
780 for (i
= 0; i
< mapnum
; i
++)
782 cur_node
.tgt_offset
= gomp_map_val (tgt
, hostaddrs
, i
);
783 /* FIXME: see above FIXME comment. */
784 devicep
->host2dev_func (devicep
->target_id
,
785 (void *) (tgt
->tgt_start
786 + i
* sizeof (void *)),
787 (void *) &cur_node
.tgt_offset
,
792 /* If the variable from "omp target enter data" map-list was already mapped,
793 tgt is not needed. Otherwise tgt will be freed by gomp_unmap_vars or
795 if (pragma_kind
== GOMP_MAP_VARS_ENTER_DATA
&& tgt
->refcount
== 0)
801 gomp_mutex_unlock (&devicep
->lock
);
806 gomp_unmap_tgt (struct target_mem_desc
*tgt
)
808 /* Deallocate on target the tgt->tgt_start .. tgt->tgt_end region. */
810 tgt
->device_descr
->free_func (tgt
->device_descr
->target_id
, tgt
->to_free
);
816 /* Decrease the refcount for a set of mapped variables, and queue asychronous
817 copies from the device back to the host after any work that has been issued.
818 Because the regions are still "live", increment an asynchronous reference
819 count to indicate that they should not be unmapped from host-side data
820 structures until the asynchronous copy has completed. */
822 attribute_hidden
void
823 gomp_copy_from_async (struct target_mem_desc
*tgt
)
825 struct gomp_device_descr
*devicep
= tgt
->device_descr
;
828 gomp_mutex_lock (&devicep
->lock
);
830 for (i
= 0; i
< tgt
->list_count
; i
++)
831 if (tgt
->list
[i
].key
== NULL
)
833 else if (tgt
->list
[i
].key
->refcount
> 1)
835 tgt
->list
[i
].key
->refcount
--;
836 tgt
->list
[i
].key
->async_refcount
++;
840 splay_tree_key k
= tgt
->list
[i
].key
;
841 if (tgt
->list
[i
].copy_from
)
842 devicep
->dev2host_func (devicep
->target_id
, (void *) k
->host_start
,
843 (void *) (k
->tgt
->tgt_start
+ k
->tgt_offset
),
844 k
->host_end
- k
->host_start
);
847 gomp_mutex_unlock (&devicep
->lock
);
850 /* Unmap variables described by TGT. If DO_COPYFROM is true, copy relevant
851 variables back from device to host: if it is false, it is assumed that this
852 has been done already, i.e. by gomp_copy_from_async above. */
854 attribute_hidden
void
855 gomp_unmap_vars (struct target_mem_desc
*tgt
, bool do_copyfrom
)
857 struct gomp_device_descr
*devicep
= tgt
->device_descr
;
859 if (tgt
->list_count
== 0)
865 gomp_mutex_lock (&devicep
->lock
);
866 if (devicep
->state
== GOMP_DEVICE_FINALIZED
)
868 gomp_mutex_unlock (&devicep
->lock
);
875 for (i
= 0; i
< tgt
->list_count
; i
++)
877 splay_tree_key k
= tgt
->list
[i
].key
;
881 bool do_unmap
= false;
882 if (k
->refcount
> 1 && k
->refcount
!= REFCOUNT_INFINITY
)
884 else if (k
->refcount
== 1)
886 if (k
->async_refcount
> 0)
895 if ((do_unmap
&& do_copyfrom
&& tgt
->list
[i
].copy_from
)
896 || tgt
->list
[i
].always_copy_from
)
897 devicep
->dev2host_func (devicep
->target_id
,
898 (void *) (k
->host_start
+ tgt
->list
[i
].offset
),
899 (void *) (k
->tgt
->tgt_start
+ k
->tgt_offset
900 + tgt
->list
[i
].offset
),
901 tgt
->list
[i
].length
);
904 splay_tree_remove (&devicep
->mem_map
, k
);
906 splay_tree_insert (&devicep
->mem_map
,
907 (splay_tree_node
) k
->link_key
);
908 if (k
->tgt
->refcount
> 1)
911 gomp_unmap_tgt (k
->tgt
);
915 if (tgt
->refcount
> 1)
918 gomp_unmap_tgt (tgt
);
920 gomp_mutex_unlock (&devicep
->lock
);
924 gomp_update (struct gomp_device_descr
*devicep
, size_t mapnum
, void **hostaddrs
,
925 size_t *sizes
, void *kinds
, bool short_mapkind
)
928 struct splay_tree_key_s cur_node
;
929 const int typemask
= short_mapkind
? 0xff : 0x7;
937 gomp_mutex_lock (&devicep
->lock
);
938 if (devicep
->state
== GOMP_DEVICE_FINALIZED
)
940 gomp_mutex_unlock (&devicep
->lock
);
944 for (i
= 0; i
< mapnum
; i
++)
947 cur_node
.host_start
= (uintptr_t) hostaddrs
[i
];
948 cur_node
.host_end
= cur_node
.host_start
+ sizes
[i
];
949 splay_tree_key n
= splay_tree_lookup (&devicep
->mem_map
, &cur_node
);
952 int kind
= get_kind (short_mapkind
, kinds
, i
);
953 if (n
->host_start
> cur_node
.host_start
954 || n
->host_end
< cur_node
.host_end
)
956 gomp_mutex_unlock (&devicep
->lock
);
957 gomp_fatal ("Trying to update [%p..%p) object when "
958 "only [%p..%p) is mapped",
959 (void *) cur_node
.host_start
,
960 (void *) cur_node
.host_end
,
961 (void *) n
->host_start
,
962 (void *) n
->host_end
);
964 if (GOMP_MAP_COPY_TO_P (kind
& typemask
))
965 devicep
->host2dev_func (devicep
->target_id
,
966 (void *) (n
->tgt
->tgt_start
968 + cur_node
.host_start
970 (void *) cur_node
.host_start
,
971 cur_node
.host_end
- cur_node
.host_start
);
972 if (GOMP_MAP_COPY_FROM_P (kind
& typemask
))
973 devicep
->dev2host_func (devicep
->target_id
,
974 (void *) cur_node
.host_start
,
975 (void *) (n
->tgt
->tgt_start
977 + cur_node
.host_start
979 cur_node
.host_end
- cur_node
.host_start
);
982 gomp_mutex_unlock (&devicep
->lock
);
985 /* Load image pointed by TARGET_DATA to the device, specified by DEVICEP.
986 And insert to splay tree the mapping between addresses from HOST_TABLE and
987 from loaded target image. We rely in the host and device compiler
988 emitting variable and functions in the same order. */
991 gomp_load_image_to_device (struct gomp_device_descr
*devicep
, unsigned version
,
992 const void *host_table
, const void *target_data
,
993 bool is_register_lock
)
995 void **host_func_table
= ((void ***) host_table
)[0];
996 void **host_funcs_end
= ((void ***) host_table
)[1];
997 void **host_var_table
= ((void ***) host_table
)[2];
998 void **host_vars_end
= ((void ***) host_table
)[3];
1000 /* The func table contains only addresses, the var table contains addresses
1001 and corresponding sizes. */
1002 int num_funcs
= host_funcs_end
- host_func_table
;
1003 int num_vars
= (host_vars_end
- host_var_table
) / 2;
1005 /* Load image to device and get target addresses for the image. */
1006 struct addr_pair
*target_table
= NULL
;
1007 int i
, num_target_entries
;
1010 = devicep
->load_image_func (devicep
->target_id
, version
,
1011 target_data
, &target_table
);
1013 if (num_target_entries
!= num_funcs
+ num_vars
)
1015 gomp_mutex_unlock (&devicep
->lock
);
1016 if (is_register_lock
)
1017 gomp_mutex_unlock (®ister_lock
);
1018 gomp_fatal ("Cannot map target functions or variables"
1019 " (expected %u, have %u)", num_funcs
+ num_vars
,
1020 num_target_entries
);
1023 /* Insert host-target address mapping into splay tree. */
1024 struct target_mem_desc
*tgt
= gomp_malloc (sizeof (*tgt
));
1025 tgt
->array
= gomp_malloc ((num_funcs
+ num_vars
) * sizeof (*tgt
->array
));
1026 tgt
->refcount
= REFCOUNT_INFINITY
;
1029 tgt
->to_free
= NULL
;
1031 tgt
->list_count
= 0;
1032 tgt
->device_descr
= devicep
;
1033 splay_tree_node array
= tgt
->array
;
1035 for (i
= 0; i
< num_funcs
; i
++)
1037 splay_tree_key k
= &array
->key
;
1038 k
->host_start
= (uintptr_t) host_func_table
[i
];
1039 k
->host_end
= k
->host_start
+ 1;
1041 k
->tgt_offset
= target_table
[i
].start
;
1042 k
->refcount
= REFCOUNT_INFINITY
;
1043 k
->async_refcount
= 0;
1046 array
->right
= NULL
;
1047 splay_tree_insert (&devicep
->mem_map
, array
);
1051 /* Most significant bit of the size in host and target tables marks
1052 "omp declare target link" variables. */
1053 const uintptr_t link_bit
= 1ULL << (sizeof (uintptr_t) * __CHAR_BIT__
- 1);
1054 const uintptr_t size_mask
= ~link_bit
;
1056 for (i
= 0; i
< num_vars
; i
++)
1058 struct addr_pair
*target_var
= &target_table
[num_funcs
+ i
];
1059 uintptr_t target_size
= target_var
->end
- target_var
->start
;
1061 if ((uintptr_t) host_var_table
[i
* 2 + 1] != target_size
)
1063 gomp_mutex_unlock (&devicep
->lock
);
1064 if (is_register_lock
)
1065 gomp_mutex_unlock (®ister_lock
);
1066 gomp_fatal ("Cannot map target variables (size mismatch)");
1069 splay_tree_key k
= &array
->key
;
1070 k
->host_start
= (uintptr_t) host_var_table
[i
* 2];
1072 = k
->host_start
+ (size_mask
& (uintptr_t) host_var_table
[i
* 2 + 1]);
1074 k
->tgt_offset
= target_var
->start
;
1075 k
->refcount
= target_size
& link_bit
? REFCOUNT_LINK
: REFCOUNT_INFINITY
;
1076 k
->async_refcount
= 0;
1079 array
->right
= NULL
;
1080 splay_tree_insert (&devicep
->mem_map
, array
);
1084 free (target_table
);
1087 /* Unload the mappings described by target_data from device DEVICE_P.
1088 The device must be locked. */
1091 gomp_unload_image_from_device (struct gomp_device_descr
*devicep
,
1093 const void *host_table
, const void *target_data
)
1095 void **host_func_table
= ((void ***) host_table
)[0];
1096 void **host_funcs_end
= ((void ***) host_table
)[1];
1097 void **host_var_table
= ((void ***) host_table
)[2];
1098 void **host_vars_end
= ((void ***) host_table
)[3];
1100 /* The func table contains only addresses, the var table contains addresses
1101 and corresponding sizes. */
1102 int num_funcs
= host_funcs_end
- host_func_table
;
1103 int num_vars
= (host_vars_end
- host_var_table
) / 2;
1105 struct splay_tree_key_s k
;
1106 splay_tree_key node
= NULL
;
1108 /* Find mapping at start of node array */
1109 if (num_funcs
|| num_vars
)
1111 k
.host_start
= (num_funcs
? (uintptr_t) host_func_table
[0]
1112 : (uintptr_t) host_var_table
[0]);
1113 k
.host_end
= k
.host_start
+ 1;
1114 node
= splay_tree_lookup (&devicep
->mem_map
, &k
);
1117 devicep
->unload_image_func (devicep
->target_id
, version
, target_data
);
1119 /* Remove mappings from splay tree. */
1121 for (i
= 0; i
< num_funcs
; i
++)
1123 k
.host_start
= (uintptr_t) host_func_table
[i
];
1124 k
.host_end
= k
.host_start
+ 1;
1125 splay_tree_remove (&devicep
->mem_map
, &k
);
1128 /* Most significant bit of the size in host and target tables marks
1129 "omp declare target link" variables. */
1130 const uintptr_t link_bit
= 1ULL << (sizeof (uintptr_t) * __CHAR_BIT__
- 1);
1131 const uintptr_t size_mask
= ~link_bit
;
1132 bool is_tgt_unmapped
= false;
1134 for (i
= 0; i
< num_vars
; i
++)
1136 k
.host_start
= (uintptr_t) host_var_table
[i
* 2];
1138 = k
.host_start
+ (size_mask
& (uintptr_t) host_var_table
[i
* 2 + 1]);
1140 if (!(link_bit
& (uintptr_t) host_var_table
[i
* 2 + 1]))
1141 splay_tree_remove (&devicep
->mem_map
, &k
);
1144 splay_tree_key n
= splay_tree_lookup (&devicep
->mem_map
, &k
);
1145 splay_tree_remove (&devicep
->mem_map
, n
);
1148 if (n
->tgt
->refcount
> 1)
1152 is_tgt_unmapped
= true;
1153 gomp_unmap_tgt (n
->tgt
);
1159 if (node
&& !is_tgt_unmapped
)
1166 /* This function should be called from every offload image while loading.
1167 It gets the descriptor of the host func and var tables HOST_TABLE, TYPE of
1168 the target, and TARGET_DATA needed by target plugin. */
1171 GOMP_offload_register_ver (unsigned version
, const void *host_table
,
1172 int target_type
, const void *target_data
)
1176 if (GOMP_VERSION_LIB (version
) > GOMP_VERSION
)
1177 gomp_fatal ("Library too old for offload (version %u < %u)",
1178 GOMP_VERSION
, GOMP_VERSION_LIB (version
));
1180 gomp_mutex_lock (®ister_lock
);
1182 /* Load image to all initialized devices. */
1183 for (i
= 0; i
< num_devices
; i
++)
1185 struct gomp_device_descr
*devicep
= &devices
[i
];
1186 gomp_mutex_lock (&devicep
->lock
);
1187 if (devicep
->type
== target_type
1188 && devicep
->state
== GOMP_DEVICE_INITIALIZED
)
1189 gomp_load_image_to_device (devicep
, version
,
1190 host_table
, target_data
, true);
1191 gomp_mutex_unlock (&devicep
->lock
);
1194 /* Insert image to array of pending images. */
1196 = gomp_realloc_unlock (offload_images
,
1197 (num_offload_images
+ 1)
1198 * sizeof (struct offload_image_descr
));
1199 offload_images
[num_offload_images
].version
= version
;
1200 offload_images
[num_offload_images
].type
= target_type
;
1201 offload_images
[num_offload_images
].host_table
= host_table
;
1202 offload_images
[num_offload_images
].target_data
= target_data
;
1204 num_offload_images
++;
1205 gomp_mutex_unlock (®ister_lock
);
1209 GOMP_offload_register (const void *host_table
, int target_type
,
1210 const void *target_data
)
1212 GOMP_offload_register_ver (0, host_table
, target_type
, target_data
);
1215 /* This function should be called from every offload image while unloading.
1216 It gets the descriptor of the host func and var tables HOST_TABLE, TYPE of
1217 the target, and TARGET_DATA needed by target plugin. */
1220 GOMP_offload_unregister_ver (unsigned version
, const void *host_table
,
1221 int target_type
, const void *target_data
)
1225 gomp_mutex_lock (®ister_lock
);
1227 /* Unload image from all initialized devices. */
1228 for (i
= 0; i
< num_devices
; i
++)
1230 struct gomp_device_descr
*devicep
= &devices
[i
];
1231 gomp_mutex_lock (&devicep
->lock
);
1232 if (devicep
->type
== target_type
1233 && devicep
->state
== GOMP_DEVICE_INITIALIZED
)
1234 gomp_unload_image_from_device (devicep
, version
,
1235 host_table
, target_data
);
1236 gomp_mutex_unlock (&devicep
->lock
);
1239 /* Remove image from array of pending images. */
1240 for (i
= 0; i
< num_offload_images
; i
++)
1241 if (offload_images
[i
].target_data
== target_data
)
1243 offload_images
[i
] = offload_images
[--num_offload_images
];
1247 gomp_mutex_unlock (®ister_lock
);
1251 GOMP_offload_unregister (const void *host_table
, int target_type
,
1252 const void *target_data
)
1254 GOMP_offload_unregister_ver (0, host_table
, target_type
, target_data
);
1257 /* This function initializes the target device, specified by DEVICEP. DEVICEP
1258 must be locked on entry, and remains locked on return. */
1260 attribute_hidden
void
1261 gomp_init_device (struct gomp_device_descr
*devicep
)
1264 devicep
->init_device_func (devicep
->target_id
);
1266 /* Load to device all images registered by the moment. */
1267 for (i
= 0; i
< num_offload_images
; i
++)
1269 struct offload_image_descr
*image
= &offload_images
[i
];
1270 if (image
->type
== devicep
->type
)
1271 gomp_load_image_to_device (devicep
, image
->version
,
1272 image
->host_table
, image
->target_data
,
1276 devicep
->state
= GOMP_DEVICE_INITIALIZED
;
1279 attribute_hidden
void
1280 gomp_unload_device (struct gomp_device_descr
*devicep
)
1282 if (devicep
->state
== GOMP_DEVICE_INITIALIZED
)
1286 /* Unload from device all images registered at the moment. */
1287 for (i
= 0; i
< num_offload_images
; i
++)
1289 struct offload_image_descr
*image
= &offload_images
[i
];
1290 if (image
->type
== devicep
->type
)
1291 gomp_unload_image_from_device (devicep
, image
->version
,
1293 image
->target_data
);
1298 /* Free address mapping tables. MM must be locked on entry, and remains locked
1301 attribute_hidden
void
1302 gomp_free_memmap (struct splay_tree_s
*mem_map
)
1304 while (mem_map
->root
)
1306 struct target_mem_desc
*tgt
= mem_map
->root
->key
.tgt
;
1308 splay_tree_remove (mem_map
, &mem_map
->root
->key
);
1314 /* Host fallback for GOMP_target{,_ext} routines. */
1317 gomp_target_fallback (void (*fn
) (void *), void **hostaddrs
)
1319 struct gomp_thread old_thr
, *thr
= gomp_thread ();
1321 memset (thr
, '\0', sizeof (*thr
));
1322 if (gomp_places_list
)
1324 thr
->place
= old_thr
.place
;
1325 thr
->ts
.place_partition_len
= gomp_places_list_len
;
1328 gomp_free_thread (thr
);
1332 /* Host fallback with firstprivate map-type handling. */
1335 gomp_target_fallback_firstprivate (void (*fn
) (void *), size_t mapnum
,
1336 void **hostaddrs
, size_t *sizes
,
1337 unsigned short *kinds
)
1339 size_t i
, tgt_align
= 0, tgt_size
= 0;
1341 for (i
= 0; i
< mapnum
; i
++)
1342 if ((kinds
[i
] & 0xff) == GOMP_MAP_FIRSTPRIVATE
)
1344 size_t align
= (size_t) 1 << (kinds
[i
] >> 8);
1345 if (tgt_align
< align
)
1347 tgt_size
= (tgt_size
+ align
- 1) & ~(align
- 1);
1348 tgt_size
+= sizes
[i
];
1352 tgt
= gomp_alloca (tgt_size
+ tgt_align
- 1);
1353 uintptr_t al
= (uintptr_t) tgt
& (tgt_align
- 1);
1355 tgt
+= tgt_align
- al
;
1357 for (i
= 0; i
< mapnum
; i
++)
1358 if ((kinds
[i
] & 0xff) == GOMP_MAP_FIRSTPRIVATE
)
1360 size_t align
= (size_t) 1 << (kinds
[i
] >> 8);
1361 tgt_size
= (tgt_size
+ align
- 1) & ~(align
- 1);
1362 memcpy (tgt
+ tgt_size
, hostaddrs
[i
], sizes
[i
]);
1363 hostaddrs
[i
] = tgt
+ tgt_size
;
1364 tgt_size
= tgt_size
+ sizes
[i
];
1367 gomp_target_fallback (fn
, hostaddrs
);
1370 /* Helper function of GOMP_target{,_ext} routines. */
1373 gomp_get_target_fn_addr (struct gomp_device_descr
*devicep
,
1374 void (*host_fn
) (void *))
1376 if (devicep
->capabilities
& GOMP_OFFLOAD_CAP_NATIVE_EXEC
)
1377 return (void *) host_fn
;
1380 gomp_mutex_lock (&devicep
->lock
);
1381 if (devicep
->state
== GOMP_DEVICE_FINALIZED
)
1383 gomp_mutex_unlock (&devicep
->lock
);
1387 struct splay_tree_key_s k
;
1388 k
.host_start
= (uintptr_t) host_fn
;
1389 k
.host_end
= k
.host_start
+ 1;
1390 splay_tree_key tgt_fn
= splay_tree_lookup (&devicep
->mem_map
, &k
);
1391 gomp_mutex_unlock (&devicep
->lock
);
1393 gomp_fatal ("Target function wasn't mapped");
1395 return (void *) tgt_fn
->tgt_offset
;
1399 /* Called when encountering a target directive. If DEVICE
1400 is GOMP_DEVICE_ICV, it means use device-var ICV. If it is
1401 GOMP_DEVICE_HOST_FALLBACK (or any value
1402 larger than last available hw device), use host fallback.
1403 FN is address of host code, UNUSED is part of the current ABI, but
1404 we're not actually using it. HOSTADDRS, SIZES and KINDS are arrays
1405 with MAPNUM entries, with addresses of the host objects,
1406 sizes of the host objects (resp. for pointer kind pointer bias
1407 and assumed sizeof (void *) size) and kinds. */
1410 GOMP_target (int device
, void (*fn
) (void *), const void *unused
,
1411 size_t mapnum
, void **hostaddrs
, size_t *sizes
,
1412 unsigned char *kinds
)
1414 struct gomp_device_descr
*devicep
= resolve_device (device
);
1418 || !(devicep
->capabilities
& GOMP_OFFLOAD_CAP_OPENMP_400
)
1419 || !(fn_addr
= gomp_get_target_fn_addr (devicep
, fn
)))
1420 return gomp_target_fallback (fn
, hostaddrs
);
1422 struct target_mem_desc
*tgt_vars
1423 = gomp_map_vars (devicep
, mapnum
, hostaddrs
, NULL
, sizes
, kinds
, false,
1424 GOMP_MAP_VARS_TARGET
);
1425 devicep
->run_func (devicep
->target_id
, fn_addr
, (void *) tgt_vars
->tgt_start
);
1426 gomp_unmap_vars (tgt_vars
, true);
1429 /* Like GOMP_target, but KINDS is 16-bit, UNUSED is no longer present,
1430 and several arguments have been added:
1431 FLAGS is a bitmask, see GOMP_TARGET_FLAG_* in gomp-constants.h.
1432 DEPEND is array of dependencies, see GOMP_task for details.
1433 NUM_TEAMS is positive if GOMP_teams will be called in the body with
1434 that value, or 1 if teams construct is not present, or 0, if
1435 teams construct does not have num_teams clause and so the choice is
1436 implementation defined, and -1 if it can't be determined on the host
1437 what value will GOMP_teams have on the device.
1438 THREAD_LIMIT similarly is positive if GOMP_teams will be called in the
1439 body with that value, or 0, if teams construct does not have thread_limit
1440 clause or the teams construct is not present, or -1 if it can't be
1441 determined on the host what value will GOMP_teams have on the device. */
1444 GOMP_target_ext (int device
, void (*fn
) (void *), size_t mapnum
,
1445 void **hostaddrs
, size_t *sizes
, unsigned short *kinds
,
1446 unsigned int flags
, void **depend
, int num_teams
,
1449 struct gomp_device_descr
*devicep
= resolve_device (device
);
1452 (void) thread_limit
;
1454 if (flags
& GOMP_TARGET_FLAG_NOWAIT
)
1456 struct gomp_thread
*thr
= gomp_thread ();
1457 /* Create a team if we don't have any around, as nowait
1458 target tasks make sense to run asynchronously even when
1459 outside of any parallel. */
1460 if (__builtin_expect (thr
->ts
.team
== NULL
, 0))
1462 struct gomp_team
*team
= gomp_new_team (1);
1463 struct gomp_task
*task
= thr
->task
;
1464 struct gomp_task_icv
*icv
= task
? &task
->icv
: &gomp_global_icv
;
1465 team
->prev_ts
= thr
->ts
;
1466 thr
->ts
.team
= team
;
1467 thr
->ts
.team_id
= 0;
1468 thr
->ts
.work_share
= &team
->work_shares
[0];
1469 thr
->ts
.last_work_share
= NULL
;
1470 #ifdef HAVE_SYNC_BUILTINS
1471 thr
->ts
.single_count
= 0;
1473 thr
->ts
.static_trip
= 0;
1474 thr
->task
= &team
->implicit_task
[0];
1475 gomp_init_task (thr
->task
, NULL
, icv
);
1481 thr
->task
= &team
->implicit_task
[0];
1484 pthread_setspecific (gomp_thread_destructor
, thr
);
1487 && !thr
->task
->final_task
)
1489 gomp_create_target_task (devicep
, fn
, mapnum
, hostaddrs
,
1490 sizes
, kinds
, flags
, depend
,
1491 GOMP_TARGET_TASK_BEFORE_MAP
);
1496 /* If there are depend clauses, but nowait is not present
1497 (or we are in a final task), block the parent task until the
1498 dependencies are resolved and then just continue with the rest
1499 of the function as if it is a merged task. */
1502 struct gomp_thread
*thr
= gomp_thread ();
1503 if (thr
->task
&& thr
->task
->depend_hash
)
1504 gomp_task_maybe_wait_for_dependencies (depend
);
1509 || !(devicep
->capabilities
& GOMP_OFFLOAD_CAP_OPENMP_400
)
1510 || !(fn_addr
= gomp_get_target_fn_addr (devicep
, fn
)))
1512 gomp_target_fallback_firstprivate (fn
, mapnum
, hostaddrs
, sizes
, kinds
);
1516 struct target_mem_desc
*tgt_vars
1517 = gomp_map_vars (devicep
, mapnum
, hostaddrs
, NULL
, sizes
, kinds
, true,
1518 GOMP_MAP_VARS_TARGET
);
1519 devicep
->run_func (devicep
->target_id
, fn_addr
, (void *) tgt_vars
->tgt_start
);
1520 gomp_unmap_vars (tgt_vars
, true);
1523 /* Host fallback for GOMP_target_data{,_ext} routines. */
1526 gomp_target_data_fallback (void)
1528 struct gomp_task_icv
*icv
= gomp_icv (false);
1529 if (icv
->target_data
)
1531 /* Even when doing a host fallback, if there are any active
1532 #pragma omp target data constructs, need to remember the
1533 new #pragma omp target data, otherwise GOMP_target_end_data
1534 would get out of sync. */
1535 struct target_mem_desc
*tgt
1536 = gomp_map_vars (NULL
, 0, NULL
, NULL
, NULL
, NULL
, false,
1537 GOMP_MAP_VARS_DATA
);
1538 tgt
->prev
= icv
->target_data
;
1539 icv
->target_data
= tgt
;
1544 GOMP_target_data (int device
, const void *unused
, size_t mapnum
,
1545 void **hostaddrs
, size_t *sizes
, unsigned char *kinds
)
1547 struct gomp_device_descr
*devicep
= resolve_device (device
);
1550 || !(devicep
->capabilities
& GOMP_OFFLOAD_CAP_OPENMP_400
))
1551 return gomp_target_data_fallback ();
1553 struct target_mem_desc
*tgt
1554 = gomp_map_vars (devicep
, mapnum
, hostaddrs
, NULL
, sizes
, kinds
, false,
1555 GOMP_MAP_VARS_DATA
);
1556 struct gomp_task_icv
*icv
= gomp_icv (true);
1557 tgt
->prev
= icv
->target_data
;
1558 icv
->target_data
= tgt
;
1562 GOMP_target_data_ext (int device
, size_t mapnum
, void **hostaddrs
,
1563 size_t *sizes
, unsigned short *kinds
)
1565 struct gomp_device_descr
*devicep
= resolve_device (device
);
1568 || !(devicep
->capabilities
& GOMP_OFFLOAD_CAP_OPENMP_400
))
1569 return gomp_target_data_fallback ();
1571 struct target_mem_desc
*tgt
1572 = gomp_map_vars (devicep
, mapnum
, hostaddrs
, NULL
, sizes
, kinds
, true,
1573 GOMP_MAP_VARS_DATA
);
1574 struct gomp_task_icv
*icv
= gomp_icv (true);
1575 tgt
->prev
= icv
->target_data
;
1576 icv
->target_data
= tgt
;
1580 GOMP_target_end_data (void)
1582 struct gomp_task_icv
*icv
= gomp_icv (false);
1583 if (icv
->target_data
)
1585 struct target_mem_desc
*tgt
= icv
->target_data
;
1586 icv
->target_data
= tgt
->prev
;
1587 gomp_unmap_vars (tgt
, true);
1592 GOMP_target_update (int device
, const void *unused
, size_t mapnum
,
1593 void **hostaddrs
, size_t *sizes
, unsigned char *kinds
)
1595 struct gomp_device_descr
*devicep
= resolve_device (device
);
1598 || !(devicep
->capabilities
& GOMP_OFFLOAD_CAP_OPENMP_400
))
1601 gomp_update (devicep
, mapnum
, hostaddrs
, sizes
, kinds
, false);
1605 GOMP_target_update_ext (int device
, size_t mapnum
, void **hostaddrs
,
1606 size_t *sizes
, unsigned short *kinds
,
1607 unsigned int flags
, void **depend
)
1609 struct gomp_device_descr
*devicep
= resolve_device (device
);
1611 /* If there are depend clauses, but nowait is not present,
1612 block the parent task until the dependencies are resolved
1613 and then just continue with the rest of the function as if it
1614 is a merged task. Until we are able to schedule task during
1615 variable mapping or unmapping, ignore nowait if depend clauses
1619 struct gomp_thread
*thr
= gomp_thread ();
1620 if (thr
->task
&& thr
->task
->depend_hash
)
1622 if ((flags
& GOMP_TARGET_FLAG_NOWAIT
)
1624 && !thr
->task
->final_task
)
1626 if (gomp_create_target_task (devicep
, (void (*) (void *)) NULL
,
1627 mapnum
, hostaddrs
, sizes
, kinds
,
1628 flags
| GOMP_TARGET_FLAG_UPDATE
,
1629 depend
, GOMP_TARGET_TASK_DATA
))
1634 struct gomp_team
*team
= thr
->ts
.team
;
1635 /* If parallel or taskgroup has been cancelled, don't start new
1638 && (gomp_team_barrier_cancelled (&team
->barrier
)
1639 || (thr
->task
->taskgroup
1640 && thr
->task
->taskgroup
->cancelled
)))
1643 gomp_task_maybe_wait_for_dependencies (depend
);
1649 || !(devicep
->capabilities
& GOMP_OFFLOAD_CAP_OPENMP_400
))
1652 struct gomp_thread
*thr
= gomp_thread ();
1653 struct gomp_team
*team
= thr
->ts
.team
;
1654 /* If parallel or taskgroup has been cancelled, don't start new tasks. */
1656 && (gomp_team_barrier_cancelled (&team
->barrier
)
1657 || (thr
->task
->taskgroup
&& thr
->task
->taskgroup
->cancelled
)))
1660 gomp_update (devicep
, mapnum
, hostaddrs
, sizes
, kinds
, true);
1664 gomp_exit_data (struct gomp_device_descr
*devicep
, size_t mapnum
,
1665 void **hostaddrs
, size_t *sizes
, unsigned short *kinds
)
1667 const int typemask
= 0xff;
1669 gomp_mutex_lock (&devicep
->lock
);
1670 if (devicep
->state
== GOMP_DEVICE_FINALIZED
)
1672 gomp_mutex_unlock (&devicep
->lock
);
1676 for (i
= 0; i
< mapnum
; i
++)
1678 struct splay_tree_key_s cur_node
;
1679 unsigned char kind
= kinds
[i
] & typemask
;
1683 case GOMP_MAP_ALWAYS_FROM
:
1684 case GOMP_MAP_DELETE
:
1685 case GOMP_MAP_RELEASE
:
1686 case GOMP_MAP_ZERO_LEN_ARRAY_SECTION
:
1687 case GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION
:
1688 cur_node
.host_start
= (uintptr_t) hostaddrs
[i
];
1689 cur_node
.host_end
= cur_node
.host_start
+ sizes
[i
];
1690 splay_tree_key k
= (kind
== GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION
1691 || kind
== GOMP_MAP_ZERO_LEN_ARRAY_SECTION
)
1692 ? gomp_map_0len_lookup (&devicep
->mem_map
, &cur_node
)
1693 : splay_tree_lookup (&devicep
->mem_map
, &cur_node
);
1697 if (k
->refcount
> 0 && k
->refcount
!= REFCOUNT_INFINITY
)
1699 if ((kind
== GOMP_MAP_DELETE
1700 || kind
== GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION
)
1701 && k
->refcount
!= REFCOUNT_INFINITY
)
1704 if ((kind
== GOMP_MAP_FROM
&& k
->refcount
== 0)
1705 || kind
== GOMP_MAP_ALWAYS_FROM
)
1706 devicep
->dev2host_func (devicep
->target_id
,
1707 (void *) cur_node
.host_start
,
1708 (void *) (k
->tgt
->tgt_start
+ k
->tgt_offset
1709 + cur_node
.host_start
1711 cur_node
.host_end
- cur_node
.host_start
);
1712 if (k
->refcount
== 0)
1714 splay_tree_remove (&devicep
->mem_map
, k
);
1716 splay_tree_insert (&devicep
->mem_map
,
1717 (splay_tree_node
) k
->link_key
);
1718 if (k
->tgt
->refcount
> 1)
1721 gomp_unmap_tgt (k
->tgt
);
1726 gomp_mutex_unlock (&devicep
->lock
);
1727 gomp_fatal ("GOMP_target_enter_exit_data unhandled kind 0x%.2x",
1732 gomp_mutex_unlock (&devicep
->lock
);
1736 GOMP_target_enter_exit_data (int device
, size_t mapnum
, void **hostaddrs
,
1737 size_t *sizes
, unsigned short *kinds
,
1738 unsigned int flags
, void **depend
)
1740 struct gomp_device_descr
*devicep
= resolve_device (device
);
1742 /* If there are depend clauses, but nowait is not present,
1743 block the parent task until the dependencies are resolved
1744 and then just continue with the rest of the function as if it
1745 is a merged task. Until we are able to schedule task during
1746 variable mapping or unmapping, ignore nowait if depend clauses
1750 struct gomp_thread
*thr
= gomp_thread ();
1751 if (thr
->task
&& thr
->task
->depend_hash
)
1753 if ((flags
& GOMP_TARGET_FLAG_NOWAIT
)
1755 && !thr
->task
->final_task
)
1757 if (gomp_create_target_task (devicep
, (void (*) (void *)) NULL
,
1758 mapnum
, hostaddrs
, sizes
, kinds
,
1760 GOMP_TARGET_TASK_DATA
))
1765 struct gomp_team
*team
= thr
->ts
.team
;
1766 /* If parallel or taskgroup has been cancelled, don't start new
1769 && (gomp_team_barrier_cancelled (&team
->barrier
)
1770 || (thr
->task
->taskgroup
1771 && thr
->task
->taskgroup
->cancelled
)))
1774 gomp_task_maybe_wait_for_dependencies (depend
);
1780 || !(devicep
->capabilities
& GOMP_OFFLOAD_CAP_OPENMP_400
))
1783 struct gomp_thread
*thr
= gomp_thread ();
1784 struct gomp_team
*team
= thr
->ts
.team
;
1785 /* If parallel or taskgroup has been cancelled, don't start new tasks. */
1787 && (gomp_team_barrier_cancelled (&team
->barrier
)
1788 || (thr
->task
->taskgroup
&& thr
->task
->taskgroup
->cancelled
)))
1792 if ((flags
& GOMP_TARGET_FLAG_EXIT_DATA
) == 0)
1793 for (i
= 0; i
< mapnum
; i
++)
1794 if ((kinds
[i
] & 0xff) == GOMP_MAP_STRUCT
)
1796 gomp_map_vars (devicep
, sizes
[i
] + 1, &hostaddrs
[i
], NULL
, &sizes
[i
],
1797 &kinds
[i
], true, GOMP_MAP_VARS_ENTER_DATA
);
1801 gomp_map_vars (devicep
, 1, &hostaddrs
[i
], NULL
, &sizes
[i
], &kinds
[i
],
1802 true, GOMP_MAP_VARS_ENTER_DATA
);
1804 gomp_exit_data (devicep
, mapnum
, hostaddrs
, sizes
, kinds
);
1808 gomp_target_task_fn (void *data
)
1810 struct gomp_target_task
*ttask
= (struct gomp_target_task
*) data
;
1811 struct gomp_device_descr
*devicep
= ttask
->devicep
;
1813 if (ttask
->fn
!= NULL
)
1817 || !(devicep
->capabilities
& GOMP_OFFLOAD_CAP_OPENMP_400
)
1818 || !(fn_addr
= gomp_get_target_fn_addr (devicep
, ttask
->fn
)))
1820 ttask
->state
= GOMP_TARGET_TASK_FALLBACK
;
1821 gomp_target_fallback_firstprivate (ttask
->fn
, ttask
->mapnum
,
1822 ttask
->hostaddrs
, ttask
->sizes
,
1827 if (ttask
->state
== GOMP_TARGET_TASK_FINISHED
)
1829 gomp_unmap_vars (ttask
->tgt
, true);
1834 = gomp_map_vars (devicep
, ttask
->mapnum
, ttask
->hostaddrs
, NULL
,
1835 ttask
->sizes
, ttask
->kinds
, true,
1836 GOMP_MAP_VARS_TARGET
);
1837 ttask
->state
= GOMP_TARGET_TASK_READY_TO_RUN
;
1839 devicep
->async_run_func (devicep
->target_id
, fn_addr
,
1840 (void *) ttask
->tgt
->tgt_start
, (void *) ttask
);
1843 else if (devicep
== NULL
1844 || !(devicep
->capabilities
& GOMP_OFFLOAD_CAP_OPENMP_400
))
1848 if (ttask
->flags
& GOMP_TARGET_FLAG_UPDATE
)
1849 gomp_update (devicep
, ttask
->mapnum
, ttask
->hostaddrs
, ttask
->sizes
,
1850 ttask
->kinds
, true);
1851 else if ((ttask
->flags
& GOMP_TARGET_FLAG_EXIT_DATA
) == 0)
1852 for (i
= 0; i
< ttask
->mapnum
; i
++)
1853 if ((ttask
->kinds
[i
] & 0xff) == GOMP_MAP_STRUCT
)
1855 gomp_map_vars (devicep
, ttask
->sizes
[i
] + 1, &ttask
->hostaddrs
[i
],
1856 NULL
, &ttask
->sizes
[i
], &ttask
->kinds
[i
], true,
1857 GOMP_MAP_VARS_ENTER_DATA
);
1858 i
+= ttask
->sizes
[i
];
1861 gomp_map_vars (devicep
, 1, &ttask
->hostaddrs
[i
], NULL
, &ttask
->sizes
[i
],
1862 &ttask
->kinds
[i
], true, GOMP_MAP_VARS_ENTER_DATA
);
1864 gomp_exit_data (devicep
, ttask
->mapnum
, ttask
->hostaddrs
, ttask
->sizes
,
1870 GOMP_teams (unsigned int num_teams
, unsigned int thread_limit
)
1874 struct gomp_task_icv
*icv
= gomp_icv (true);
1875 icv
->thread_limit_var
1876 = thread_limit
> INT_MAX
? UINT_MAX
: thread_limit
;
1882 omp_target_alloc (size_t size
, int device_num
)
1884 if (device_num
== GOMP_DEVICE_HOST_FALLBACK
)
1885 return malloc (size
);
1890 struct gomp_device_descr
*devicep
= resolve_device (device_num
);
1891 if (devicep
== NULL
)
1894 if (!(devicep
->capabilities
& GOMP_OFFLOAD_CAP_OPENMP_400
))
1895 return malloc (size
);
1897 gomp_mutex_lock (&devicep
->lock
);
1898 void *ret
= devicep
->alloc_func (devicep
->target_id
, size
);
1899 gomp_mutex_unlock (&devicep
->lock
);
1904 omp_target_free (void *device_ptr
, int device_num
)
1906 if (device_ptr
== NULL
)
1909 if (device_num
== GOMP_DEVICE_HOST_FALLBACK
)
1918 struct gomp_device_descr
*devicep
= resolve_device (device_num
);
1919 if (devicep
== NULL
)
1922 if (!(devicep
->capabilities
& GOMP_OFFLOAD_CAP_OPENMP_400
))
1928 gomp_mutex_lock (&devicep
->lock
);
1929 devicep
->free_func (devicep
->target_id
, device_ptr
);
1930 gomp_mutex_unlock (&devicep
->lock
);
1934 omp_target_is_present (void *ptr
, int device_num
)
1939 if (device_num
== GOMP_DEVICE_HOST_FALLBACK
)
1945 struct gomp_device_descr
*devicep
= resolve_device (device_num
);
1946 if (devicep
== NULL
)
1949 if (!(devicep
->capabilities
& GOMP_OFFLOAD_CAP_OPENMP_400
))
1952 gomp_mutex_lock (&devicep
->lock
);
1953 struct splay_tree_s
*mem_map
= &devicep
->mem_map
;
1954 struct splay_tree_key_s cur_node
;
1956 cur_node
.host_start
= (uintptr_t) ptr
;
1957 cur_node
.host_end
= cur_node
.host_start
;
1958 splay_tree_key n
= gomp_map_0len_lookup (mem_map
, &cur_node
);
1959 int ret
= n
!= NULL
;
1960 gomp_mutex_unlock (&devicep
->lock
);
1965 omp_target_memcpy (void *dst
, void *src
, size_t length
, size_t dst_offset
,
1966 size_t src_offset
, int dst_device_num
, int src_device_num
)
1968 struct gomp_device_descr
*dst_devicep
= NULL
, *src_devicep
= NULL
;
1970 if (dst_device_num
!= GOMP_DEVICE_HOST_FALLBACK
)
1972 if (dst_device_num
< 0)
1975 dst_devicep
= resolve_device (dst_device_num
);
1976 if (dst_devicep
== NULL
)
1979 if (!(dst_devicep
->capabilities
& GOMP_OFFLOAD_CAP_OPENMP_400
))
1982 if (src_device_num
!= GOMP_DEVICE_HOST_FALLBACK
)
1984 if (src_device_num
< 0)
1987 src_devicep
= resolve_device (src_device_num
);
1988 if (src_devicep
== NULL
)
1991 if (!(src_devicep
->capabilities
& GOMP_OFFLOAD_CAP_OPENMP_400
))
1994 if (src_devicep
== NULL
&& dst_devicep
== NULL
)
1996 memcpy ((char *) dst
+ dst_offset
, (char *) src
+ src_offset
, length
);
1999 if (src_devicep
== NULL
)
2001 gomp_mutex_lock (&dst_devicep
->lock
);
2002 dst_devicep
->host2dev_func (dst_devicep
->target_id
,
2003 (char *) dst
+ dst_offset
,
2004 (char *) src
+ src_offset
, length
);
2005 gomp_mutex_unlock (&dst_devicep
->lock
);
2008 if (dst_devicep
== NULL
)
2010 gomp_mutex_lock (&src_devicep
->lock
);
2011 src_devicep
->dev2host_func (src_devicep
->target_id
,
2012 (char *) dst
+ dst_offset
,
2013 (char *) src
+ src_offset
, length
);
2014 gomp_mutex_unlock (&src_devicep
->lock
);
2017 if (src_devicep
== dst_devicep
)
2019 gomp_mutex_lock (&src_devicep
->lock
);
2020 src_devicep
->dev2dev_func (src_devicep
->target_id
,
2021 (char *) dst
+ dst_offset
,
2022 (char *) src
+ src_offset
, length
);
2023 gomp_mutex_unlock (&src_devicep
->lock
);
2030 omp_target_memcpy_rect_worker (void *dst
, void *src
, size_t element_size
,
2031 int num_dims
, const size_t *volume
,
2032 const size_t *dst_offsets
,
2033 const size_t *src_offsets
,
2034 const size_t *dst_dimensions
,
2035 const size_t *src_dimensions
,
2036 struct gomp_device_descr
*dst_devicep
,
2037 struct gomp_device_descr
*src_devicep
)
2039 size_t dst_slice
= element_size
;
2040 size_t src_slice
= element_size
;
2041 size_t j
, dst_off
, src_off
, length
;
2046 if (__builtin_mul_overflow (element_size
, volume
[0], &length
)
2047 || __builtin_mul_overflow (element_size
, dst_offsets
[0], &dst_off
)
2048 || __builtin_mul_overflow (element_size
, src_offsets
[0], &src_off
))
2050 if (dst_devicep
== NULL
&& src_devicep
== NULL
)
2051 memcpy ((char *) dst
+ dst_off
, (char *) src
+ src_off
, length
);
2052 else if (src_devicep
== NULL
)
2053 dst_devicep
->host2dev_func (dst_devicep
->target_id
,
2054 (char *) dst
+ dst_off
,
2055 (char *) src
+ src_off
, length
);
2056 else if (dst_devicep
== NULL
)
2057 src_devicep
->dev2host_func (src_devicep
->target_id
,
2058 (char *) dst
+ dst_off
,
2059 (char *) src
+ src_off
, length
);
2060 else if (src_devicep
== dst_devicep
)
2061 src_devicep
->dev2dev_func (src_devicep
->target_id
,
2062 (char *) dst
+ dst_off
,
2063 (char *) src
+ src_off
, length
);
2069 /* FIXME: it would be nice to have some plugin function to handle
2070 num_dims == 2 and num_dims == 3 more efficiently. Larger ones can
2071 be handled in the generic recursion below, and for host-host it
2072 should be used even for any num_dims >= 2. */
2074 for (i
= 1; i
< num_dims
; i
++)
2075 if (__builtin_mul_overflow (dst_slice
, dst_dimensions
[i
], &dst_slice
)
2076 || __builtin_mul_overflow (src_slice
, src_dimensions
[i
], &src_slice
))
2078 if (__builtin_mul_overflow (dst_slice
, dst_offsets
[0], &dst_off
)
2079 || __builtin_mul_overflow (src_slice
, src_offsets
[0], &src_off
))
2081 for (j
= 0; j
< volume
[0]; j
++)
2083 ret
= omp_target_memcpy_rect_worker ((char *) dst
+ dst_off
,
2084 (char *) src
+ src_off
,
2085 element_size
, num_dims
- 1,
2086 volume
+ 1, dst_offsets
+ 1,
2087 src_offsets
+ 1, dst_dimensions
+ 1,
2088 src_dimensions
+ 1, dst_devicep
,
2092 dst_off
+= dst_slice
;
2093 src_off
+= src_slice
;
2099 omp_target_memcpy_rect (void *dst
, void *src
, size_t element_size
,
2100 int num_dims
, const size_t *volume
,
2101 const size_t *dst_offsets
,
2102 const size_t *src_offsets
,
2103 const size_t *dst_dimensions
,
2104 const size_t *src_dimensions
,
2105 int dst_device_num
, int src_device_num
)
2107 struct gomp_device_descr
*dst_devicep
= NULL
, *src_devicep
= NULL
;
2112 if (dst_device_num
!= GOMP_DEVICE_HOST_FALLBACK
)
2114 if (dst_device_num
< 0)
2117 dst_devicep
= resolve_device (dst_device_num
);
2118 if (dst_devicep
== NULL
)
2121 if (!(dst_devicep
->capabilities
& GOMP_OFFLOAD_CAP_OPENMP_400
))
2124 if (src_device_num
!= GOMP_DEVICE_HOST_FALLBACK
)
2126 if (src_device_num
< 0)
2129 src_devicep
= resolve_device (src_device_num
);
2130 if (src_devicep
== NULL
)
2133 if (!(src_devicep
->capabilities
& GOMP_OFFLOAD_CAP_OPENMP_400
))
2137 if (src_devicep
!= NULL
&& dst_devicep
!= NULL
&& src_devicep
!= dst_devicep
)
2141 gomp_mutex_lock (&src_devicep
->lock
);
2142 else if (dst_devicep
)
2143 gomp_mutex_lock (&dst_devicep
->lock
);
2144 int ret
= omp_target_memcpy_rect_worker (dst
, src
, element_size
, num_dims
,
2145 volume
, dst_offsets
, src_offsets
,
2146 dst_dimensions
, src_dimensions
,
2147 dst_devicep
, src_devicep
);
2149 gomp_mutex_unlock (&src_devicep
->lock
);
2150 else if (dst_devicep
)
2151 gomp_mutex_unlock (&dst_devicep
->lock
);
2156 omp_target_associate_ptr (void *host_ptr
, void *device_ptr
, size_t size
,
2157 size_t device_offset
, int device_num
)
2159 if (device_num
== GOMP_DEVICE_HOST_FALLBACK
)
2165 struct gomp_device_descr
*devicep
= resolve_device (device_num
);
2166 if (devicep
== NULL
)
2169 if (!(devicep
->capabilities
& GOMP_OFFLOAD_CAP_OPENMP_400
))
2172 gomp_mutex_lock (&devicep
->lock
);
2174 struct splay_tree_s
*mem_map
= &devicep
->mem_map
;
2175 struct splay_tree_key_s cur_node
;
2178 cur_node
.host_start
= (uintptr_t) host_ptr
;
2179 cur_node
.host_end
= cur_node
.host_start
+ size
;
2180 splay_tree_key n
= gomp_map_lookup (mem_map
, &cur_node
);
2183 if (n
->tgt
->tgt_start
+ n
->tgt_offset
2184 == (uintptr_t) device_ptr
+ device_offset
2185 && n
->host_start
<= cur_node
.host_start
2186 && n
->host_end
>= cur_node
.host_end
)
2191 struct target_mem_desc
*tgt
= gomp_malloc (sizeof (*tgt
));
2192 tgt
->array
= gomp_malloc (sizeof (*tgt
->array
));
2196 tgt
->to_free
= NULL
;
2198 tgt
->list_count
= 0;
2199 tgt
->device_descr
= devicep
;
2200 splay_tree_node array
= tgt
->array
;
2201 splay_tree_key k
= &array
->key
;
2202 k
->host_start
= cur_node
.host_start
;
2203 k
->host_end
= cur_node
.host_end
;
2205 k
->tgt_offset
= (uintptr_t) device_ptr
+ device_offset
;
2206 k
->refcount
= REFCOUNT_INFINITY
;
2207 k
->async_refcount
= 0;
2209 array
->right
= NULL
;
2210 splay_tree_insert (&devicep
->mem_map
, array
);
2213 gomp_mutex_unlock (&devicep
->lock
);
2218 omp_target_disassociate_ptr (void *ptr
, int device_num
)
2220 if (device_num
== GOMP_DEVICE_HOST_FALLBACK
)
2226 struct gomp_device_descr
*devicep
= resolve_device (device_num
);
2227 if (devicep
== NULL
)
2230 if (!(devicep
->capabilities
& GOMP_OFFLOAD_CAP_OPENMP_400
))
2233 gomp_mutex_lock (&devicep
->lock
);
2235 struct splay_tree_s
*mem_map
= &devicep
->mem_map
;
2236 struct splay_tree_key_s cur_node
;
2239 cur_node
.host_start
= (uintptr_t) ptr
;
2240 cur_node
.host_end
= cur_node
.host_start
;
2241 splay_tree_key n
= gomp_map_lookup (mem_map
, &cur_node
);
2243 && n
->host_start
== cur_node
.host_start
2244 && n
->refcount
== REFCOUNT_INFINITY
2245 && n
->tgt
->tgt_start
== 0
2246 && n
->tgt
->to_free
== NULL
2247 && n
->tgt
->refcount
== 1
2248 && n
->tgt
->list_count
== 0)
2250 splay_tree_remove (&devicep
->mem_map
, n
);
2251 gomp_unmap_tgt (n
->tgt
);
2255 gomp_mutex_unlock (&devicep
->lock
);
2259 #ifdef PLUGIN_SUPPORT
2261 /* This function tries to load a plugin for DEVICE. Name of plugin is passed
2263 The handles of the found functions are stored in the corresponding fields
2264 of DEVICE. The function returns TRUE on success and FALSE otherwise. */
2267 gomp_load_plugin_for_device (struct gomp_device_descr
*device
,
2268 const char *plugin_name
)
2270 const char *err
= NULL
, *last_missing
= NULL
;
2272 void *plugin_handle
= dlopen (plugin_name
, RTLD_LAZY
);
2276 /* Check if all required functions are available in the plugin and store
2277 their handlers. None of the symbols can legitimately be NULL,
2278 so we don't need to check dlerror all the time. */
2280 if (!(device->f##_func = dlsym (plugin_handle, "GOMP_OFFLOAD_" #f))) \
2282 /* Similar, but missing functions are not an error. Return false if
2283 failed, true otherwise. */
2284 #define DLSYM_OPT(f, n) \
2285 ((device->f##_func = dlsym (plugin_handle, "GOMP_OFFLOAD_" #n)) \
2286 || (last_missing = #n, 0))
2289 if (device
->version_func () != GOMP_VERSION
)
2291 err
= "plugin version mismatch";
2298 DLSYM (get_num_devices
);
2299 DLSYM (init_device
);
2300 DLSYM (fini_device
);
2302 DLSYM (unload_image
);
2307 device
->capabilities
= device
->get_caps_func ();
2308 if (device
->capabilities
& GOMP_OFFLOAD_CAP_OPENMP_400
)
2314 if (device
->capabilities
& GOMP_OFFLOAD_CAP_OPENACC_200
)
2316 if (!DLSYM_OPT (openacc
.exec
, openacc_parallel
)
2317 || !DLSYM_OPT (openacc
.register_async_cleanup
,
2318 openacc_register_async_cleanup
)
2319 || !DLSYM_OPT (openacc
.async_test
, openacc_async_test
)
2320 || !DLSYM_OPT (openacc
.async_test_all
, openacc_async_test_all
)
2321 || !DLSYM_OPT (openacc
.async_wait
, openacc_async_wait
)
2322 || !DLSYM_OPT (openacc
.async_wait_async
, openacc_async_wait_async
)
2323 || !DLSYM_OPT (openacc
.async_wait_all
, openacc_async_wait_all
)
2324 || !DLSYM_OPT (openacc
.async_wait_all_async
,
2325 openacc_async_wait_all_async
)
2326 || !DLSYM_OPT (openacc
.async_set_async
, openacc_async_set_async
)
2327 || !DLSYM_OPT (openacc
.create_thread_data
,
2328 openacc_create_thread_data
)
2329 || !DLSYM_OPT (openacc
.destroy_thread_data
,
2330 openacc_destroy_thread_data
))
2332 /* Require all the OpenACC handlers if we have
2333 GOMP_OFFLOAD_CAP_OPENACC_200. */
2334 err
= "plugin missing OpenACC handler function";
2339 cuda
+= DLSYM_OPT (openacc
.cuda
.get_current_device
,
2340 openacc_get_current_cuda_device
);
2341 cuda
+= DLSYM_OPT (openacc
.cuda
.get_current_context
,
2342 openacc_get_current_cuda_context
);
2343 cuda
+= DLSYM_OPT (openacc
.cuda
.get_stream
, openacc_get_cuda_stream
);
2344 cuda
+= DLSYM_OPT (openacc
.cuda
.set_stream
, openacc_set_cuda_stream
);
2345 if (cuda
&& cuda
!= 4)
2347 /* Make sure all the CUDA functions are there if any of them are. */
2348 err
= "plugin missing OpenACC CUDA handler function";
2360 gomp_error ("while loading %s: %s", plugin_name
, err
);
2362 gomp_error ("missing function was %s", last_missing
);
2364 dlclose (plugin_handle
);
2369 /* This function finalizes all initialized devices. */
2372 gomp_target_fini (void)
2375 for (i
= 0; i
< num_devices
; i
++)
2377 struct gomp_device_descr
*devicep
= &devices
[i
];
2378 gomp_mutex_lock (&devicep
->lock
);
2379 if (devicep
->state
== GOMP_DEVICE_INITIALIZED
)
2381 devicep
->fini_device_func (devicep
->target_id
);
2382 devicep
->state
= GOMP_DEVICE_FINALIZED
;
2384 gomp_mutex_unlock (&devicep
->lock
);
2388 /* This function initializes the runtime needed for offloading.
2389 It parses the list of offload targets and tries to load the plugins for
2390 these targets. On return, the variables NUM_DEVICES and NUM_DEVICES_OPENMP
2391 will be set, and the array DEVICES initialized, containing descriptors for
2392 corresponding devices, first the GOMP_OFFLOAD_CAP_OPENMP_400 ones, follows
2396 gomp_target_init (void)
2398 const char *prefix
="libgomp-plugin-";
2399 const char *suffix
= SONAME_SUFFIX (1);
2400 const char *cur
, *next
;
2402 int i
, new_num_devices
;
2407 cur
= OFFLOAD_TARGETS
;
2411 struct gomp_device_descr current_device
;
2413 next
= strchr (cur
, ',');
2415 plugin_name
= (char *) malloc (1 + (next
? next
- cur
: strlen (cur
))
2416 + strlen (prefix
) + strlen (suffix
));
2423 strcpy (plugin_name
, prefix
);
2424 strncat (plugin_name
, cur
, next
? next
- cur
: strlen (cur
));
2425 strcat (plugin_name
, suffix
);
2427 if (gomp_load_plugin_for_device (¤t_device
, plugin_name
))
2429 new_num_devices
= current_device
.get_num_devices_func ();
2430 if (new_num_devices
>= 1)
2432 /* Augment DEVICES and NUM_DEVICES. */
2434 devices
= realloc (devices
, (num_devices
+ new_num_devices
)
2435 * sizeof (struct gomp_device_descr
));
2443 current_device
.name
= current_device
.get_name_func ();
2444 /* current_device.capabilities has already been set. */
2445 current_device
.type
= current_device
.get_type_func ();
2446 current_device
.mem_map
.root
= NULL
;
2447 current_device
.state
= GOMP_DEVICE_UNINITIALIZED
;
2448 current_device
.openacc
.data_environ
= NULL
;
2449 for (i
= 0; i
< new_num_devices
; i
++)
2451 current_device
.target_id
= i
;
2452 devices
[num_devices
] = current_device
;
2453 gomp_mutex_init (&devices
[num_devices
].lock
);
2464 /* In DEVICES, sort the GOMP_OFFLOAD_CAP_OPENMP_400 ones first, and set
2465 NUM_DEVICES_OPENMP. */
2466 struct gomp_device_descr
*devices_s
2467 = malloc (num_devices
* sizeof (struct gomp_device_descr
));
2474 num_devices_openmp
= 0;
2475 for (i
= 0; i
< num_devices
; i
++)
2476 if (devices
[i
].capabilities
& GOMP_OFFLOAD_CAP_OPENMP_400
)
2477 devices_s
[num_devices_openmp
++] = devices
[i
];
2478 int num_devices_after_openmp
= num_devices_openmp
;
2479 for (i
= 0; i
< num_devices
; i
++)
2480 if (!(devices
[i
].capabilities
& GOMP_OFFLOAD_CAP_OPENMP_400
))
2481 devices_s
[num_devices_after_openmp
++] = devices
[i
];
2483 devices
= devices_s
;
2485 for (i
= 0; i
< num_devices
; i
++)
2487 /* The 'devices' array can be moved (by the realloc call) until we have
2488 found all the plugins, so registering with the OpenACC runtime (which
2489 takes a copy of the pointer argument) must be delayed until now. */
2490 if (devices
[i
].capabilities
& GOMP_OFFLOAD_CAP_OPENACC_200
)
2491 goacc_register (&devices
[i
]);
2494 if (atexit (gomp_target_fini
) != 0)
2495 gomp_fatal ("atexit failed");
2498 #else /* PLUGIN_SUPPORT */
2499 /* If dlfcn.h is unavailable we always fallback to host execution.
2500 GOMP_target* routines are just stubs for this case. */
2502 gomp_target_init (void)
2505 #endif /* PLUGIN_SUPPORT */