1 /* Copyright (C) 2013-2020 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. */
29 #include "oacc-plugin.h"
31 #include "gomp-constants.h"
35 #ifdef HAVE_INTTYPES_H
36 # include <inttypes.h> /* For PRIu64. */
44 #include "plugin-suffix.h"
47 #define FIELD_TGT_EMPTY (~(size_t) 0)
49 static void gomp_target_init (void);
51 /* The whole initialization code for offloading plugins is only run one. */
52 static pthread_once_t gomp_is_initialized
= PTHREAD_ONCE_INIT
;
54 /* Mutex for offload image registration. */
55 static gomp_mutex_t register_lock
;
57 /* This structure describes an offload image.
58 It contains type of the target device, pointer to host table descriptor, and
59 pointer to target data. */
60 struct offload_image_descr
{
62 enum offload_target_type type
;
63 const void *host_table
;
64 const void *target_data
;
67 /* Array of descriptors of offload images. */
68 static struct offload_image_descr
*offload_images
;
70 /* Total number of offload images. */
71 static int num_offload_images
;
73 /* Array of descriptors for all available devices. */
74 static struct gomp_device_descr
*devices
;
76 /* Total number of available devices. */
77 static int num_devices
;
79 /* Number of GOMP_OFFLOAD_CAP_OPENMP_400 devices. */
80 static int num_devices_openmp
;
82 /* Similar to gomp_realloc, but release register_lock before gomp_fatal. */
85 gomp_realloc_unlock (void *old
, size_t size
)
87 void *ret
= realloc (old
, size
);
90 gomp_mutex_unlock (®ister_lock
);
91 gomp_fatal ("Out of memory allocating %lu bytes", (unsigned long) size
);
97 gomp_init_targets_once (void)
99 (void) pthread_once (&gomp_is_initialized
, gomp_target_init
);
103 gomp_get_num_devices (void)
105 gomp_init_targets_once ();
106 return num_devices_openmp
;
109 static struct gomp_device_descr
*
110 resolve_device (int device_id
)
112 if (device_id
== GOMP_DEVICE_ICV
)
114 struct gomp_task_icv
*icv
= gomp_icv (false);
115 device_id
= icv
->default_device_var
;
118 if (device_id
< 0 || device_id
>= gomp_get_num_devices ())
121 gomp_mutex_lock (&devices
[device_id
].lock
);
122 if (devices
[device_id
].state
== GOMP_DEVICE_UNINITIALIZED
)
123 gomp_init_device (&devices
[device_id
]);
124 else if (devices
[device_id
].state
== GOMP_DEVICE_FINALIZED
)
126 gomp_mutex_unlock (&devices
[device_id
].lock
);
129 gomp_mutex_unlock (&devices
[device_id
].lock
);
131 return &devices
[device_id
];
135 static inline splay_tree_key
136 gomp_map_lookup (splay_tree mem_map
, splay_tree_key key
)
138 if (key
->host_start
!= key
->host_end
)
139 return splay_tree_lookup (mem_map
, key
);
142 splay_tree_key n
= splay_tree_lookup (mem_map
, key
);
147 n
= splay_tree_lookup (mem_map
, key
);
151 return splay_tree_lookup (mem_map
, key
);
154 static inline splay_tree_key
155 gomp_map_0len_lookup (splay_tree mem_map
, splay_tree_key key
)
157 if (key
->host_start
!= key
->host_end
)
158 return splay_tree_lookup (mem_map
, key
);
161 splay_tree_key n
= splay_tree_lookup (mem_map
, key
);
167 gomp_device_copy (struct gomp_device_descr
*devicep
,
168 bool (*copy_func
) (int, void *, const void *, size_t),
169 const char *dst
, void *dstaddr
,
170 const char *src
, const void *srcaddr
,
173 if (!copy_func (devicep
->target_id
, dstaddr
, srcaddr
, size
))
175 gomp_mutex_unlock (&devicep
->lock
);
176 gomp_fatal ("Copying of %s object [%p..%p) to %s object [%p..%p) failed",
177 src
, srcaddr
, srcaddr
+ size
, dst
, dstaddr
, dstaddr
+ size
);
182 goacc_device_copy_async (struct gomp_device_descr
*devicep
,
183 bool (*copy_func
) (int, void *, const void *, size_t,
184 struct goacc_asyncqueue
*),
185 const char *dst
, void *dstaddr
,
186 const char *src
, const void *srcaddr
,
187 size_t size
, struct goacc_asyncqueue
*aq
)
189 if (!copy_func (devicep
->target_id
, dstaddr
, srcaddr
, size
, aq
))
191 gomp_mutex_unlock (&devicep
->lock
);
192 gomp_fatal ("Copying of %s object [%p..%p) to %s object [%p..%p) failed",
193 src
, srcaddr
, srcaddr
+ size
, dst
, dstaddr
, dstaddr
+ size
);
197 /* Infrastructure for coalescing adjacent or nearly adjacent (in device addresses)
198 host to device memory transfers. */
200 struct gomp_coalesce_chunk
202 /* The starting and ending point of a coalesced chunk of memory. */
206 struct gomp_coalesce_buf
208 /* Buffer into which gomp_copy_host2dev will memcpy data and from which
209 it will be copied to the device. */
211 struct target_mem_desc
*tgt
;
212 /* Array with offsets, chunks[i].start is the starting offset and
213 chunks[i].end ending offset relative to tgt->tgt_start device address
214 of chunks which are to be copied to buf and later copied to device. */
215 struct gomp_coalesce_chunk
*chunks
;
216 /* Number of chunks in chunks array, or -1 if coalesce buffering should not
219 /* During construction of chunks array, how many memory regions are within
220 the last chunk. If there is just one memory region for a chunk, we copy
221 it directly to device rather than going through buf. */
225 /* Maximum size of memory region considered for coalescing. Larger copies
226 are performed directly. */
227 #define MAX_COALESCE_BUF_SIZE (32 * 1024)
229 /* Maximum size of a gap in between regions to consider them being copied
230 within the same chunk. All the device offsets considered are within
231 newly allocated device memory, so it isn't fatal if we copy some padding
232 in between from host to device. The gaps come either from alignment
233 padding or from memory regions which are not supposed to be copied from
234 host to device (e.g. map(alloc:), map(from:) etc.). */
235 #define MAX_COALESCE_BUF_GAP (4 * 1024)
237 /* Add region with device tgt_start relative offset and length to CBUF. */
240 gomp_coalesce_buf_add (struct gomp_coalesce_buf
*cbuf
, size_t start
, size_t len
)
242 if (len
> MAX_COALESCE_BUF_SIZE
|| len
== 0)
246 if (cbuf
->chunk_cnt
< 0)
248 if (start
< cbuf
->chunks
[cbuf
->chunk_cnt
- 1].end
)
250 cbuf
->chunk_cnt
= -1;
253 if (start
< cbuf
->chunks
[cbuf
->chunk_cnt
- 1].end
+ MAX_COALESCE_BUF_GAP
)
255 cbuf
->chunks
[cbuf
->chunk_cnt
- 1].end
= start
+ len
;
259 /* If the last chunk is only used by one mapping, discard it,
260 as it will be one host to device copy anyway and
261 memcpying it around will only waste cycles. */
262 if (cbuf
->use_cnt
== 1)
265 cbuf
->chunks
[cbuf
->chunk_cnt
].start
= start
;
266 cbuf
->chunks
[cbuf
->chunk_cnt
].end
= start
+ len
;
271 /* Return true for mapping kinds which need to copy data from the
272 host to device for regions that weren't previously mapped. */
275 gomp_to_device_kind_p (int kind
)
281 case GOMP_MAP_FORCE_ALLOC
:
282 case GOMP_MAP_FORCE_FROM
:
283 case GOMP_MAP_ALWAYS_FROM
:
290 attribute_hidden
void
291 gomp_copy_host2dev (struct gomp_device_descr
*devicep
,
292 struct goacc_asyncqueue
*aq
,
293 void *d
, const void *h
, size_t sz
,
294 struct gomp_coalesce_buf
*cbuf
)
298 uintptr_t doff
= (uintptr_t) d
- cbuf
->tgt
->tgt_start
;
299 if (doff
< cbuf
->chunks
[cbuf
->chunk_cnt
- 1].end
)
302 long last
= cbuf
->chunk_cnt
- 1;
303 while (first
<= last
)
305 long middle
= (first
+ last
) >> 1;
306 if (cbuf
->chunks
[middle
].end
<= doff
)
308 else if (cbuf
->chunks
[middle
].start
<= doff
)
310 if (doff
+ sz
> cbuf
->chunks
[middle
].end
)
311 gomp_fatal ("internal libgomp cbuf error");
312 memcpy ((char *) cbuf
->buf
+ (doff
- cbuf
->chunks
[0].start
),
321 if (__builtin_expect (aq
!= NULL
, 0))
322 goacc_device_copy_async (devicep
, devicep
->openacc
.async
.host2dev_func
,
323 "dev", d
, "host", h
, sz
, aq
);
325 gomp_device_copy (devicep
, devicep
->host2dev_func
, "dev", d
, "host", h
, sz
);
328 attribute_hidden
void
329 gomp_copy_dev2host (struct gomp_device_descr
*devicep
,
330 struct goacc_asyncqueue
*aq
,
331 void *h
, const void *d
, size_t sz
)
333 if (__builtin_expect (aq
!= NULL
, 0))
334 goacc_device_copy_async (devicep
, devicep
->openacc
.async
.dev2host_func
,
335 "host", h
, "dev", d
, sz
, aq
);
337 gomp_device_copy (devicep
, devicep
->dev2host_func
, "host", h
, "dev", d
, sz
);
341 gomp_free_device_memory (struct gomp_device_descr
*devicep
, void *devptr
)
343 if (!devicep
->free_func (devicep
->target_id
, devptr
))
345 gomp_mutex_unlock (&devicep
->lock
);
346 gomp_fatal ("error in freeing device memory block at %p", devptr
);
350 /* Handle the case where gomp_map_lookup, splay_tree_lookup or
351 gomp_map_0len_lookup found oldn for newn.
352 Helper function of gomp_map_vars. */
355 gomp_map_vars_existing (struct gomp_device_descr
*devicep
,
356 struct goacc_asyncqueue
*aq
, splay_tree_key oldn
,
357 splay_tree_key newn
, struct target_var_desc
*tgt_var
,
358 unsigned char kind
, struct gomp_coalesce_buf
*cbuf
)
360 assert (kind
!= GOMP_MAP_ATTACH
);
363 tgt_var
->copy_from
= GOMP_MAP_COPY_FROM_P (kind
);
364 tgt_var
->always_copy_from
= GOMP_MAP_ALWAYS_FROM_P (kind
);
365 tgt_var
->is_attach
= false;
366 tgt_var
->offset
= newn
->host_start
- oldn
->host_start
;
367 tgt_var
->length
= newn
->host_end
- newn
->host_start
;
369 if ((kind
& GOMP_MAP_FLAG_FORCE
)
370 || oldn
->host_start
> newn
->host_start
371 || oldn
->host_end
< newn
->host_end
)
373 gomp_mutex_unlock (&devicep
->lock
);
374 gomp_fatal ("Trying to map into device [%p..%p) object when "
375 "[%p..%p) is already mapped",
376 (void *) newn
->host_start
, (void *) newn
->host_end
,
377 (void *) oldn
->host_start
, (void *) oldn
->host_end
);
380 if (GOMP_MAP_ALWAYS_TO_P (kind
))
381 gomp_copy_host2dev (devicep
, aq
,
382 (void *) (oldn
->tgt
->tgt_start
+ oldn
->tgt_offset
383 + newn
->host_start
- oldn
->host_start
),
384 (void *) newn
->host_start
,
385 newn
->host_end
- newn
->host_start
, cbuf
);
387 if (oldn
->refcount
!= REFCOUNT_INFINITY
)
392 get_kind (bool short_mapkind
, void *kinds
, int idx
)
394 return short_mapkind
? ((unsigned short *) kinds
)[idx
]
395 : ((unsigned char *) kinds
)[idx
];
399 gomp_map_pointer (struct target_mem_desc
*tgt
, struct goacc_asyncqueue
*aq
,
400 uintptr_t host_ptr
, uintptr_t target_offset
, uintptr_t bias
,
401 struct gomp_coalesce_buf
*cbuf
)
403 struct gomp_device_descr
*devicep
= tgt
->device_descr
;
404 struct splay_tree_s
*mem_map
= &devicep
->mem_map
;
405 struct splay_tree_key_s cur_node
;
407 cur_node
.host_start
= host_ptr
;
408 if (cur_node
.host_start
== (uintptr_t) NULL
)
410 cur_node
.tgt_offset
= (uintptr_t) NULL
;
411 gomp_copy_host2dev (devicep
, aq
,
412 (void *) (tgt
->tgt_start
+ target_offset
),
413 (void *) &cur_node
.tgt_offset
,
414 sizeof (void *), cbuf
);
417 /* Add bias to the pointer value. */
418 cur_node
.host_start
+= bias
;
419 cur_node
.host_end
= cur_node
.host_start
;
420 splay_tree_key n
= gomp_map_lookup (mem_map
, &cur_node
);
423 gomp_mutex_unlock (&devicep
->lock
);
424 gomp_fatal ("Pointer target of array section wasn't mapped");
426 cur_node
.host_start
-= n
->host_start
;
428 = n
->tgt
->tgt_start
+ n
->tgt_offset
+ cur_node
.host_start
;
429 /* At this point tgt_offset is target address of the
430 array section. Now subtract bias to get what we want
431 to initialize the pointer with. */
432 cur_node
.tgt_offset
-= bias
;
433 gomp_copy_host2dev (devicep
, aq
, (void *) (tgt
->tgt_start
+ target_offset
),
434 (void *) &cur_node
.tgt_offset
, sizeof (void *), cbuf
);
438 gomp_map_fields_existing (struct target_mem_desc
*tgt
,
439 struct goacc_asyncqueue
*aq
, splay_tree_key n
,
440 size_t first
, size_t i
, void **hostaddrs
,
441 size_t *sizes
, void *kinds
,
442 struct gomp_coalesce_buf
*cbuf
)
444 struct gomp_device_descr
*devicep
= tgt
->device_descr
;
445 struct splay_tree_s
*mem_map
= &devicep
->mem_map
;
446 struct splay_tree_key_s cur_node
;
448 const bool short_mapkind
= true;
449 const int typemask
= short_mapkind
? 0xff : 0x7;
451 cur_node
.host_start
= (uintptr_t) hostaddrs
[i
];
452 cur_node
.host_end
= cur_node
.host_start
+ sizes
[i
];
453 splay_tree_key n2
= splay_tree_lookup (mem_map
, &cur_node
);
454 kind
= get_kind (short_mapkind
, kinds
, i
);
457 && n2
->host_start
- n
->host_start
== n2
->tgt_offset
- n
->tgt_offset
)
459 gomp_map_vars_existing (devicep
, aq
, n2
, &cur_node
,
460 &tgt
->list
[i
], kind
& typemask
, cbuf
);
465 if (cur_node
.host_start
> (uintptr_t) hostaddrs
[first
- 1])
467 cur_node
.host_start
--;
468 n2
= splay_tree_lookup (mem_map
, &cur_node
);
469 cur_node
.host_start
++;
472 && n2
->host_start
- n
->host_start
473 == n2
->tgt_offset
- n
->tgt_offset
)
475 gomp_map_vars_existing (devicep
, aq
, n2
, &cur_node
,
476 &tgt
->list
[i
], kind
& typemask
, cbuf
);
481 n2
= splay_tree_lookup (mem_map
, &cur_node
);
485 && n2
->host_start
- n
->host_start
== n2
->tgt_offset
- n
->tgt_offset
)
487 gomp_map_vars_existing (devicep
, aq
, n2
, &cur_node
, &tgt
->list
[i
],
488 kind
& typemask
, cbuf
);
492 gomp_mutex_unlock (&devicep
->lock
);
493 gomp_fatal ("Trying to map into device [%p..%p) structure element when "
494 "other mapped elements from the same structure weren't mapped "
495 "together with it", (void *) cur_node
.host_start
,
496 (void *) cur_node
.host_end
);
499 attribute_hidden
void
500 gomp_attach_pointer (struct gomp_device_descr
*devicep
,
501 struct goacc_asyncqueue
*aq
, splay_tree mem_map
,
502 splay_tree_key n
, uintptr_t attach_to
, size_t bias
,
503 struct gomp_coalesce_buf
*cbufp
)
505 struct splay_tree_key_s s
;
510 gomp_mutex_unlock (&devicep
->lock
);
511 gomp_fatal ("enclosing struct not mapped for attach");
514 size
= (n
->host_end
- n
->host_start
+ sizeof (void *) - 1) / sizeof (void *);
515 /* We might have a pointer in a packed struct: however we cannot have more
516 than one such pointer in each pointer-sized portion of the struct, so
518 idx
= (attach_to
- n
->host_start
) / sizeof (void *);
521 n
->aux
= gomp_malloc_cleared (sizeof (struct splay_tree_aux
));
523 if (!n
->aux
->attach_count
)
525 = gomp_malloc_cleared (sizeof (*n
->aux
->attach_count
) * size
);
527 if (n
->aux
->attach_count
[idx
] < UINTPTR_MAX
)
528 n
->aux
->attach_count
[idx
]++;
531 gomp_mutex_unlock (&devicep
->lock
);
532 gomp_fatal ("attach count overflow");
535 if (n
->aux
->attach_count
[idx
] == 1)
537 uintptr_t devptr
= n
->tgt
->tgt_start
+ n
->tgt_offset
+ attach_to
539 uintptr_t target
= (uintptr_t) *(void **) attach_to
;
543 if ((void *) target
== NULL
)
545 gomp_mutex_unlock (&devicep
->lock
);
546 gomp_fatal ("attempt to attach null pointer");
549 s
.host_start
= target
+ bias
;
550 s
.host_end
= s
.host_start
+ 1;
551 tn
= splay_tree_lookup (mem_map
, &s
);
555 gomp_mutex_unlock (&devicep
->lock
);
556 gomp_fatal ("pointer target not mapped for attach");
559 data
= tn
->tgt
->tgt_start
+ tn
->tgt_offset
+ target
- tn
->host_start
;
562 "%s: attaching host %p, target %p (struct base %p) to %p\n",
563 __FUNCTION__
, (void *) attach_to
, (void *) devptr
,
564 (void *) (n
->tgt
->tgt_start
+ n
->tgt_offset
), (void *) data
);
566 gomp_copy_host2dev (devicep
, aq
, (void *) devptr
, (void *) &data
,
567 sizeof (void *), cbufp
);
570 gomp_debug (1, "%s: attach count for %p -> %u\n", __FUNCTION__
,
571 (void *) attach_to
, (int) n
->aux
->attach_count
[idx
]);
574 attribute_hidden
void
575 gomp_detach_pointer (struct gomp_device_descr
*devicep
,
576 struct goacc_asyncqueue
*aq
, splay_tree_key n
,
577 uintptr_t detach_from
, bool finalize
,
578 struct gomp_coalesce_buf
*cbufp
)
584 gomp_mutex_unlock (&devicep
->lock
);
585 gomp_fatal ("enclosing struct not mapped for detach");
588 idx
= (detach_from
- n
->host_start
) / sizeof (void *);
590 if (!n
->aux
|| !n
->aux
->attach_count
)
592 gomp_mutex_unlock (&devicep
->lock
);
593 gomp_fatal ("no attachment counters for struct");
597 n
->aux
->attach_count
[idx
] = 1;
599 if (n
->aux
->attach_count
[idx
] == 0)
601 gomp_mutex_unlock (&devicep
->lock
);
602 gomp_fatal ("attach count underflow");
605 n
->aux
->attach_count
[idx
]--;
607 if (n
->aux
->attach_count
[idx
] == 0)
609 uintptr_t devptr
= n
->tgt
->tgt_start
+ n
->tgt_offset
+ detach_from
611 uintptr_t target
= (uintptr_t) *(void **) detach_from
;
614 "%s: detaching host %p, target %p (struct base %p) to %p\n",
615 __FUNCTION__
, (void *) detach_from
, (void *) devptr
,
616 (void *) (n
->tgt
->tgt_start
+ n
->tgt_offset
),
619 gomp_copy_host2dev (devicep
, aq
, (void *) devptr
, (void *) &target
,
620 sizeof (void *), cbufp
);
623 gomp_debug (1, "%s: attach count for %p -> %u\n", __FUNCTION__
,
624 (void *) detach_from
, (int) n
->aux
->attach_count
[idx
]);
627 attribute_hidden
uintptr_t
628 gomp_map_val (struct target_mem_desc
*tgt
, void **hostaddrs
, size_t i
)
630 if (tgt
->list
[i
].key
!= NULL
)
631 return tgt
->list
[i
].key
->tgt
->tgt_start
632 + tgt
->list
[i
].key
->tgt_offset
633 + tgt
->list
[i
].offset
;
635 switch (tgt
->list
[i
].offset
)
638 return (uintptr_t) hostaddrs
[i
];
644 return tgt
->list
[i
+ 1].key
->tgt
->tgt_start
645 + tgt
->list
[i
+ 1].key
->tgt_offset
646 + tgt
->list
[i
+ 1].offset
647 + (uintptr_t) hostaddrs
[i
]
648 - (uintptr_t) hostaddrs
[i
+ 1];
651 return tgt
->tgt_start
+ tgt
->list
[i
].offset
;
655 static inline __attribute__((always_inline
)) struct target_mem_desc
*
656 gomp_map_vars_internal (struct gomp_device_descr
*devicep
,
657 struct goacc_asyncqueue
*aq
, size_t mapnum
,
658 void **hostaddrs
, void **devaddrs
, size_t *sizes
,
659 void *kinds
, bool short_mapkind
,
660 enum gomp_map_vars_kind pragma_kind
)
662 size_t i
, tgt_align
, tgt_size
, not_found_cnt
= 0;
663 bool has_firstprivate
= false;
664 const int rshift
= short_mapkind
? 8 : 3;
665 const int typemask
= short_mapkind
? 0xff : 0x7;
666 struct splay_tree_s
*mem_map
= &devicep
->mem_map
;
667 struct splay_tree_key_s cur_node
;
668 struct target_mem_desc
*tgt
669 = gomp_malloc (sizeof (*tgt
) + sizeof (tgt
->list
[0]) * mapnum
);
670 tgt
->list_count
= mapnum
;
671 tgt
->refcount
= pragma_kind
== GOMP_MAP_VARS_ENTER_DATA
? 0 : 1;
672 tgt
->device_descr
= devicep
;
674 struct gomp_coalesce_buf cbuf
, *cbufp
= NULL
;
683 tgt_align
= sizeof (void *);
689 if (mapnum
> 1 || pragma_kind
== GOMP_MAP_VARS_TARGET
)
691 size_t chunks_size
= (mapnum
+ 1) * sizeof (struct gomp_coalesce_chunk
);
692 cbuf
.chunks
= (struct gomp_coalesce_chunk
*) gomp_alloca (chunks_size
);
695 if (pragma_kind
== GOMP_MAP_VARS_TARGET
)
697 size_t align
= 4 * sizeof (void *);
699 tgt_size
= mapnum
* sizeof (void *);
701 cbuf
.use_cnt
= 1 + (mapnum
> 1);
702 cbuf
.chunks
[0].start
= 0;
703 cbuf
.chunks
[0].end
= tgt_size
;
706 gomp_mutex_lock (&devicep
->lock
);
707 if (devicep
->state
== GOMP_DEVICE_FINALIZED
)
709 gomp_mutex_unlock (&devicep
->lock
);
714 for (i
= 0; i
< mapnum
; i
++)
716 int kind
= get_kind (short_mapkind
, kinds
, i
);
717 if (hostaddrs
[i
] == NULL
718 || (kind
& typemask
) == GOMP_MAP_FIRSTPRIVATE_INT
)
720 tgt
->list
[i
].key
= NULL
;
721 tgt
->list
[i
].offset
= OFFSET_INLINED
;
724 else if ((kind
& typemask
) == GOMP_MAP_USE_DEVICE_PTR
725 || (kind
& typemask
) == GOMP_MAP_USE_DEVICE_PTR_IF_PRESENT
)
727 tgt
->list
[i
].key
= NULL
;
730 /* In OpenMP < 5.0 and OpenACC the mapping has to be done
731 on a separate construct prior to using use_device_{addr,ptr}.
732 In OpenMP 5.0, map directives need to be ordered by the
733 middle-end before the use_device_* clauses. If
734 !not_found_cnt, all mappings requested (if any) are already
735 mapped, so use_device_{addr,ptr} can be resolved right away.
736 Otherwise, if not_found_cnt, gomp_map_lookup might fail
737 now but would succeed after performing the mappings in the
738 following loop. We can't defer this always to the second
739 loop, because it is not even invoked when !not_found_cnt
740 after the first loop. */
741 cur_node
.host_start
= (uintptr_t) hostaddrs
[i
];
742 cur_node
.host_end
= cur_node
.host_start
;
743 splay_tree_key n
= gomp_map_lookup (mem_map
, &cur_node
);
746 cur_node
.host_start
-= n
->host_start
;
748 = (void *) (n
->tgt
->tgt_start
+ n
->tgt_offset
749 + cur_node
.host_start
);
751 else if ((kind
& typemask
) == GOMP_MAP_USE_DEVICE_PTR
)
753 gomp_mutex_unlock (&devicep
->lock
);
754 gomp_fatal ("use_device_ptr pointer wasn't mapped");
756 else if ((kind
& typemask
) == GOMP_MAP_USE_DEVICE_PTR_IF_PRESENT
)
757 /* If not present, continue using the host address. */
760 __builtin_unreachable ();
761 tgt
->list
[i
].offset
= OFFSET_INLINED
;
764 tgt
->list
[i
].offset
= 0;
767 else if ((kind
& typemask
) == GOMP_MAP_STRUCT
)
769 size_t first
= i
+ 1;
770 size_t last
= i
+ sizes
[i
];
771 cur_node
.host_start
= (uintptr_t) hostaddrs
[i
];
772 cur_node
.host_end
= (uintptr_t) hostaddrs
[last
]
774 tgt
->list
[i
].key
= NULL
;
775 tgt
->list
[i
].offset
= OFFSET_STRUCT
;
776 splay_tree_key n
= splay_tree_lookup (mem_map
, &cur_node
);
779 size_t align
= (size_t) 1 << (kind
>> rshift
);
780 if (tgt_align
< align
)
782 tgt_size
-= (uintptr_t) hostaddrs
[first
] - cur_node
.host_start
;
783 tgt_size
= (tgt_size
+ align
- 1) & ~(align
- 1);
784 tgt_size
+= cur_node
.host_end
- cur_node
.host_start
;
785 not_found_cnt
+= last
- i
;
786 for (i
= first
; i
<= last
; i
++)
788 tgt
->list
[i
].key
= NULL
;
789 if (gomp_to_device_kind_p (get_kind (short_mapkind
, kinds
, i
)
791 gomp_coalesce_buf_add (&cbuf
,
792 tgt_size
- cur_node
.host_end
793 + (uintptr_t) hostaddrs
[i
],
799 for (i
= first
; i
<= last
; i
++)
800 gomp_map_fields_existing (tgt
, aq
, n
, first
, i
, hostaddrs
,
805 else if ((kind
& typemask
) == GOMP_MAP_ALWAYS_POINTER
)
807 tgt
->list
[i
].key
= NULL
;
808 tgt
->list
[i
].offset
= OFFSET_POINTER
;
809 has_firstprivate
= true;
812 else if ((kind
& typemask
) == GOMP_MAP_ATTACH
)
814 tgt
->list
[i
].key
= NULL
;
815 has_firstprivate
= true;
818 cur_node
.host_start
= (uintptr_t) hostaddrs
[i
];
819 if (!GOMP_MAP_POINTER_P (kind
& typemask
))
820 cur_node
.host_end
= cur_node
.host_start
+ sizes
[i
];
822 cur_node
.host_end
= cur_node
.host_start
+ sizeof (void *);
823 if ((kind
& typemask
) == GOMP_MAP_FIRSTPRIVATE
)
825 tgt
->list
[i
].key
= NULL
;
827 size_t align
= (size_t) 1 << (kind
>> rshift
);
828 if (tgt_align
< align
)
830 tgt_size
= (tgt_size
+ align
- 1) & ~(align
- 1);
831 gomp_coalesce_buf_add (&cbuf
, tgt_size
,
832 cur_node
.host_end
- cur_node
.host_start
);
833 tgt_size
+= cur_node
.host_end
- cur_node
.host_start
;
834 has_firstprivate
= true;
838 if ((kind
& typemask
) == GOMP_MAP_ZERO_LEN_ARRAY_SECTION
)
840 n
= gomp_map_0len_lookup (mem_map
, &cur_node
);
843 tgt
->list
[i
].key
= NULL
;
844 tgt
->list
[i
].offset
= OFFSET_POINTER
;
849 n
= splay_tree_lookup (mem_map
, &cur_node
);
850 if (n
&& n
->refcount
!= REFCOUNT_LINK
)
851 gomp_map_vars_existing (devicep
, aq
, n
, &cur_node
, &tgt
->list
[i
],
852 kind
& typemask
, NULL
);
855 tgt
->list
[i
].key
= NULL
;
857 if ((kind
& typemask
) == GOMP_MAP_IF_PRESENT
)
859 /* Not present, hence, skip entry - including its MAP_POINTER,
861 tgt
->list
[i
].offset
= OFFSET_POINTER
;
863 && ((typemask
& get_kind (short_mapkind
, kinds
, i
+ 1))
864 == GOMP_MAP_POINTER
))
867 tgt
->list
[i
].key
= NULL
;
868 tgt
->list
[i
].offset
= 0;
872 size_t align
= (size_t) 1 << (kind
>> rshift
);
874 if (tgt_align
< align
)
876 tgt_size
= (tgt_size
+ align
- 1) & ~(align
- 1);
877 if (gomp_to_device_kind_p (kind
& typemask
))
878 gomp_coalesce_buf_add (&cbuf
, tgt_size
,
879 cur_node
.host_end
- cur_node
.host_start
);
880 tgt_size
+= cur_node
.host_end
- cur_node
.host_start
;
881 if ((kind
& typemask
) == GOMP_MAP_TO_PSET
)
884 for (j
= i
+ 1; j
< mapnum
; j
++)
885 if (!GOMP_MAP_POINTER_P (get_kind (short_mapkind
, kinds
, j
)
888 else if ((uintptr_t) hostaddrs
[j
] < cur_node
.host_start
889 || ((uintptr_t) hostaddrs
[j
] + sizeof (void *)
890 > cur_node
.host_end
))
894 tgt
->list
[j
].key
= NULL
;
905 gomp_mutex_unlock (&devicep
->lock
);
906 gomp_fatal ("unexpected aggregation");
908 tgt
->to_free
= devaddrs
[0];
909 tgt
->tgt_start
= (uintptr_t) tgt
->to_free
;
910 tgt
->tgt_end
= tgt
->tgt_start
+ sizes
[0];
912 else if (not_found_cnt
|| pragma_kind
== GOMP_MAP_VARS_TARGET
)
914 /* Allocate tgt_align aligned tgt_size block of memory. */
915 /* FIXME: Perhaps change interface to allocate properly aligned
917 tgt
->to_free
= devicep
->alloc_func (devicep
->target_id
,
918 tgt_size
+ tgt_align
- 1);
921 gomp_mutex_unlock (&devicep
->lock
);
922 gomp_fatal ("device memory allocation fail");
925 tgt
->tgt_start
= (uintptr_t) tgt
->to_free
;
926 tgt
->tgt_start
= (tgt
->tgt_start
+ tgt_align
- 1) & ~(tgt_align
- 1);
927 tgt
->tgt_end
= tgt
->tgt_start
+ tgt_size
;
929 if (cbuf
.use_cnt
== 1)
931 if (cbuf
.chunk_cnt
> 0)
934 = malloc (cbuf
.chunks
[cbuf
.chunk_cnt
- 1].end
- cbuf
.chunks
[0].start
);
950 if (pragma_kind
== GOMP_MAP_VARS_TARGET
)
951 tgt_size
= mapnum
* sizeof (void *);
954 if (not_found_cnt
|| has_firstprivate
)
957 tgt
->array
= gomp_malloc (not_found_cnt
* sizeof (*tgt
->array
));
958 splay_tree_node array
= tgt
->array
;
959 size_t j
, field_tgt_offset
= 0, field_tgt_clear
= ~(size_t) 0;
960 uintptr_t field_tgt_base
= 0;
962 for (i
= 0; i
< mapnum
; i
++)
963 if (tgt
->list
[i
].key
== NULL
)
965 int kind
= get_kind (short_mapkind
, kinds
, i
);
966 if (hostaddrs
[i
] == NULL
)
968 switch (kind
& typemask
)
970 size_t align
, len
, first
, last
;
972 case GOMP_MAP_FIRSTPRIVATE
:
973 align
= (size_t) 1 << (kind
>> rshift
);
974 tgt_size
= (tgt_size
+ align
- 1) & ~(align
- 1);
975 tgt
->list
[i
].offset
= tgt_size
;
977 gomp_copy_host2dev (devicep
, aq
,
978 (void *) (tgt
->tgt_start
+ tgt_size
),
979 (void *) hostaddrs
[i
], len
, cbufp
);
982 case GOMP_MAP_FIRSTPRIVATE_INT
:
983 case GOMP_MAP_ZERO_LEN_ARRAY_SECTION
:
985 case GOMP_MAP_USE_DEVICE_PTR_IF_PRESENT
:
986 /* The OpenACC 'host_data' construct only allows 'use_device'
987 "mapping" clauses, so in the first loop, 'not_found_cnt'
988 must always have been zero, so all OpenACC 'use_device'
989 clauses have already been handled. (We can only easily test
990 'use_device' with 'if_present' clause here.) */
991 assert (tgt
->list
[i
].offset
== OFFSET_INLINED
);
992 /* Nevertheless, FALLTHRU to the normal handling, to keep the
993 code conceptually simple, similar to the first loop. */
994 case GOMP_MAP_USE_DEVICE_PTR
:
995 if (tgt
->list
[i
].offset
== 0)
997 cur_node
.host_start
= (uintptr_t) hostaddrs
[i
];
998 cur_node
.host_end
= cur_node
.host_start
;
999 n
= gomp_map_lookup (mem_map
, &cur_node
);
1002 cur_node
.host_start
-= n
->host_start
;
1004 = (void *) (n
->tgt
->tgt_start
+ n
->tgt_offset
1005 + cur_node
.host_start
);
1007 else if ((kind
& typemask
) == GOMP_MAP_USE_DEVICE_PTR
)
1009 gomp_mutex_unlock (&devicep
->lock
);
1010 gomp_fatal ("use_device_ptr pointer wasn't mapped");
1012 else if ((kind
& typemask
)
1013 == GOMP_MAP_USE_DEVICE_PTR_IF_PRESENT
)
1014 /* If not present, continue using the host address. */
1017 __builtin_unreachable ();
1018 tgt
->list
[i
].offset
= OFFSET_INLINED
;
1021 case GOMP_MAP_STRUCT
:
1023 last
= i
+ sizes
[i
];
1024 cur_node
.host_start
= (uintptr_t) hostaddrs
[i
];
1025 cur_node
.host_end
= (uintptr_t) hostaddrs
[last
]
1027 if (tgt
->list
[first
].key
!= NULL
)
1029 n
= splay_tree_lookup (mem_map
, &cur_node
);
1032 size_t align
= (size_t) 1 << (kind
>> rshift
);
1033 tgt_size
-= (uintptr_t) hostaddrs
[first
]
1034 - (uintptr_t) hostaddrs
[i
];
1035 tgt_size
= (tgt_size
+ align
- 1) & ~(align
- 1);
1036 tgt_size
+= (uintptr_t) hostaddrs
[first
]
1037 - (uintptr_t) hostaddrs
[i
];
1038 field_tgt_base
= (uintptr_t) hostaddrs
[first
];
1039 field_tgt_offset
= tgt_size
;
1040 field_tgt_clear
= last
;
1041 tgt_size
+= cur_node
.host_end
1042 - (uintptr_t) hostaddrs
[first
];
1045 for (i
= first
; i
<= last
; i
++)
1046 gomp_map_fields_existing (tgt
, aq
, n
, first
, i
, hostaddrs
,
1047 sizes
, kinds
, cbufp
);
1050 case GOMP_MAP_ALWAYS_POINTER
:
1051 cur_node
.host_start
= (uintptr_t) hostaddrs
[i
];
1052 cur_node
.host_end
= cur_node
.host_start
+ sizeof (void *);
1053 n
= splay_tree_lookup (mem_map
, &cur_node
);
1055 || n
->host_start
> cur_node
.host_start
1056 || n
->host_end
< cur_node
.host_end
)
1058 gomp_mutex_unlock (&devicep
->lock
);
1059 gomp_fatal ("always pointer not mapped");
1061 if ((get_kind (short_mapkind
, kinds
, i
- 1) & typemask
)
1062 != GOMP_MAP_ALWAYS_POINTER
)
1063 cur_node
.tgt_offset
= gomp_map_val (tgt
, hostaddrs
, i
- 1);
1064 if (cur_node
.tgt_offset
)
1065 cur_node
.tgt_offset
-= sizes
[i
];
1066 gomp_copy_host2dev (devicep
, aq
,
1067 (void *) (n
->tgt
->tgt_start
1069 + cur_node
.host_start
1071 (void *) &cur_node
.tgt_offset
,
1072 sizeof (void *), cbufp
);
1073 cur_node
.tgt_offset
= n
->tgt
->tgt_start
+ n
->tgt_offset
1074 + cur_node
.host_start
- n
->host_start
;
1076 case GOMP_MAP_IF_PRESENT
:
1077 /* Not present - otherwise handled above. Skip over its
1078 MAP_POINTER as well. */
1080 && ((typemask
& get_kind (short_mapkind
, kinds
, i
+ 1))
1081 == GOMP_MAP_POINTER
))
1084 case GOMP_MAP_ATTACH
:
1086 cur_node
.host_start
= (uintptr_t) hostaddrs
[i
];
1087 cur_node
.host_end
= cur_node
.host_start
+ sizeof (void *);
1088 splay_tree_key n
= splay_tree_lookup (mem_map
, &cur_node
);
1091 tgt
->list
[i
].key
= n
;
1092 tgt
->list
[i
].offset
= cur_node
.host_start
- n
->host_start
;
1093 tgt
->list
[i
].length
= n
->host_end
- n
->host_start
;
1094 tgt
->list
[i
].copy_from
= false;
1095 tgt
->list
[i
].always_copy_from
= false;
1096 tgt
->list
[i
].is_attach
= true;
1097 /* OpenACC 'attach'/'detach' doesn't affect
1098 structured/dynamic reference counts ('n->refcount',
1099 'n->dynamic_refcount'). */
1103 gomp_mutex_unlock (&devicep
->lock
);
1104 gomp_fatal ("outer struct not mapped for attach");
1106 gomp_attach_pointer (devicep
, aq
, mem_map
, n
,
1107 (uintptr_t) hostaddrs
[i
], sizes
[i
],
1114 splay_tree_key k
= &array
->key
;
1115 k
->host_start
= (uintptr_t) hostaddrs
[i
];
1116 if (!GOMP_MAP_POINTER_P (kind
& typemask
))
1117 k
->host_end
= k
->host_start
+ sizes
[i
];
1119 k
->host_end
= k
->host_start
+ sizeof (void *);
1120 splay_tree_key n
= splay_tree_lookup (mem_map
, k
);
1121 if (n
&& n
->refcount
!= REFCOUNT_LINK
)
1122 gomp_map_vars_existing (devicep
, aq
, n
, k
, &tgt
->list
[i
],
1123 kind
& typemask
, cbufp
);
1127 if (n
&& n
->refcount
== REFCOUNT_LINK
)
1129 /* Replace target address of the pointer with target address
1130 of mapped object in the splay tree. */
1131 splay_tree_remove (mem_map
, n
);
1133 = gomp_malloc_cleared (sizeof (struct splay_tree_aux
));
1134 k
->aux
->link_key
= n
;
1136 size_t align
= (size_t) 1 << (kind
>> rshift
);
1137 tgt
->list
[i
].key
= k
;
1139 if (field_tgt_clear
!= FIELD_TGT_EMPTY
)
1141 k
->tgt_offset
= k
->host_start
- field_tgt_base
1143 if (i
== field_tgt_clear
)
1144 field_tgt_clear
= FIELD_TGT_EMPTY
;
1148 tgt_size
= (tgt_size
+ align
- 1) & ~(align
- 1);
1149 k
->tgt_offset
= tgt_size
;
1150 tgt_size
+= k
->host_end
- k
->host_start
;
1152 tgt
->list
[i
].copy_from
= GOMP_MAP_COPY_FROM_P (kind
& typemask
);
1153 tgt
->list
[i
].always_copy_from
1154 = GOMP_MAP_ALWAYS_FROM_P (kind
& typemask
);
1155 tgt
->list
[i
].is_attach
= false;
1156 tgt
->list
[i
].offset
= 0;
1157 tgt
->list
[i
].length
= k
->host_end
- k
->host_start
;
1159 k
->dynamic_refcount
= 0;
1162 array
->right
= NULL
;
1163 splay_tree_insert (mem_map
, array
);
1164 switch (kind
& typemask
)
1166 case GOMP_MAP_ALLOC
:
1168 case GOMP_MAP_FORCE_ALLOC
:
1169 case GOMP_MAP_FORCE_FROM
:
1170 case GOMP_MAP_ALWAYS_FROM
:
1173 case GOMP_MAP_TOFROM
:
1174 case GOMP_MAP_FORCE_TO
:
1175 case GOMP_MAP_FORCE_TOFROM
:
1176 case GOMP_MAP_ALWAYS_TO
:
1177 case GOMP_MAP_ALWAYS_TOFROM
:
1178 gomp_copy_host2dev (devicep
, aq
,
1179 (void *) (tgt
->tgt_start
1181 (void *) k
->host_start
,
1182 k
->host_end
- k
->host_start
, cbufp
);
1184 case GOMP_MAP_POINTER
:
1185 gomp_map_pointer (tgt
, aq
,
1186 (uintptr_t) *(void **) k
->host_start
,
1187 k
->tgt_offset
, sizes
[i
], cbufp
);
1189 case GOMP_MAP_TO_PSET
:
1190 gomp_copy_host2dev (devicep
, aq
,
1191 (void *) (tgt
->tgt_start
1193 (void *) k
->host_start
,
1194 k
->host_end
- k
->host_start
, cbufp
);
1196 for (j
= i
+ 1; j
< mapnum
; j
++)
1197 if (!GOMP_MAP_POINTER_P (get_kind (short_mapkind
, kinds
,
1201 else if ((uintptr_t) hostaddrs
[j
] < k
->host_start
1202 || ((uintptr_t) hostaddrs
[j
] + sizeof (void *)
1207 tgt
->list
[j
].key
= k
;
1208 tgt
->list
[j
].copy_from
= false;
1209 tgt
->list
[j
].always_copy_from
= false;
1210 tgt
->list
[j
].is_attach
= false;
1211 if (k
->refcount
!= REFCOUNT_INFINITY
)
1213 gomp_map_pointer (tgt
, aq
,
1214 (uintptr_t) *(void **) hostaddrs
[j
],
1216 + ((uintptr_t) hostaddrs
[j
]
1222 case GOMP_MAP_FORCE_PRESENT
:
1224 /* We already looked up the memory region above and it
1226 size_t size
= k
->host_end
- k
->host_start
;
1227 gomp_mutex_unlock (&devicep
->lock
);
1228 #ifdef HAVE_INTTYPES_H
1229 gomp_fatal ("present clause: !acc_is_present (%p, "
1230 "%"PRIu64
" (0x%"PRIx64
"))",
1231 (void *) k
->host_start
,
1232 (uint64_t) size
, (uint64_t) size
);
1234 gomp_fatal ("present clause: !acc_is_present (%p, "
1235 "%lu (0x%lx))", (void *) k
->host_start
,
1236 (unsigned long) size
, (unsigned long) size
);
1240 case GOMP_MAP_FORCE_DEVICEPTR
:
1241 assert (k
->host_end
- k
->host_start
== sizeof (void *));
1242 gomp_copy_host2dev (devicep
, aq
,
1243 (void *) (tgt
->tgt_start
1245 (void *) k
->host_start
,
1246 sizeof (void *), cbufp
);
1249 gomp_mutex_unlock (&devicep
->lock
);
1250 gomp_fatal ("%s: unhandled kind 0x%.2x", __FUNCTION__
,
1254 if (k
->aux
&& k
->aux
->link_key
)
1256 /* Set link pointer on target to the device address of the
1258 void *tgt_addr
= (void *) (tgt
->tgt_start
+ k
->tgt_offset
);
1259 /* We intentionally do not use coalescing here, as it's not
1260 data allocated by the current call to this function. */
1261 gomp_copy_host2dev (devicep
, aq
, (void *) n
->tgt_offset
,
1262 &tgt_addr
, sizeof (void *), NULL
);
1269 if (pragma_kind
== GOMP_MAP_VARS_TARGET
)
1271 for (i
= 0; i
< mapnum
; i
++)
1273 cur_node
.tgt_offset
= gomp_map_val (tgt
, hostaddrs
, i
);
1274 gomp_copy_host2dev (devicep
, aq
,
1275 (void *) (tgt
->tgt_start
+ i
* sizeof (void *)),
1276 (void *) &cur_node
.tgt_offset
, sizeof (void *),
1284 for (c
= 0; c
< cbuf
.chunk_cnt
; ++c
)
1285 gomp_copy_host2dev (devicep
, aq
,
1286 (void *) (tgt
->tgt_start
+ cbuf
.chunks
[c
].start
),
1287 (char *) cbuf
.buf
+ (cbuf
.chunks
[c
].start
1288 - cbuf
.chunks
[0].start
),
1289 cbuf
.chunks
[c
].end
- cbuf
.chunks
[c
].start
, NULL
);
1295 /* If the variable from "omp target enter data" map-list was already mapped,
1296 tgt is not needed. Otherwise tgt will be freed by gomp_unmap_vars or
1298 if (pragma_kind
== GOMP_MAP_VARS_ENTER_DATA
&& tgt
->refcount
== 0)
1304 gomp_mutex_unlock (&devicep
->lock
);
1308 attribute_hidden
struct target_mem_desc
*
1309 gomp_map_vars (struct gomp_device_descr
*devicep
, size_t mapnum
,
1310 void **hostaddrs
, void **devaddrs
, size_t *sizes
, void *kinds
,
1311 bool short_mapkind
, enum gomp_map_vars_kind pragma_kind
)
1313 return gomp_map_vars_internal (devicep
, NULL
, mapnum
, hostaddrs
, devaddrs
,
1314 sizes
, kinds
, short_mapkind
, pragma_kind
);
1317 attribute_hidden
struct target_mem_desc
*
1318 gomp_map_vars_async (struct gomp_device_descr
*devicep
,
1319 struct goacc_asyncqueue
*aq
, size_t mapnum
,
1320 void **hostaddrs
, void **devaddrs
, size_t *sizes
,
1321 void *kinds
, bool short_mapkind
,
1322 enum gomp_map_vars_kind pragma_kind
)
1324 return gomp_map_vars_internal (devicep
, aq
, mapnum
, hostaddrs
, devaddrs
,
1325 sizes
, kinds
, short_mapkind
, pragma_kind
);
1329 gomp_unmap_tgt (struct target_mem_desc
*tgt
)
1331 /* Deallocate on target the tgt->tgt_start .. tgt->tgt_end region. */
1333 gomp_free_device_memory (tgt
->device_descr
, tgt
->to_free
);
1340 gomp_unref_tgt (void *ptr
)
1342 bool is_tgt_unmapped
= false;
1344 struct target_mem_desc
*tgt
= (struct target_mem_desc
*) ptr
;
1346 if (tgt
->refcount
> 1)
1350 gomp_unmap_tgt (tgt
);
1351 is_tgt_unmapped
= true;
1354 return is_tgt_unmapped
;
1358 gomp_unref_tgt_void (void *ptr
)
1360 (void) gomp_unref_tgt (ptr
);
1363 static inline __attribute__((always_inline
)) bool
1364 gomp_remove_var_internal (struct gomp_device_descr
*devicep
, splay_tree_key k
,
1365 struct goacc_asyncqueue
*aq
)
1367 bool is_tgt_unmapped
= false;
1368 splay_tree_remove (&devicep
->mem_map
, k
);
1371 if (k
->aux
->link_key
)
1372 splay_tree_insert (&devicep
->mem_map
,
1373 (splay_tree_node
) k
->aux
->link_key
);
1374 if (k
->aux
->attach_count
)
1375 free (k
->aux
->attach_count
);
1380 devicep
->openacc
.async
.queue_callback_func (aq
, gomp_unref_tgt_void
,
1383 is_tgt_unmapped
= gomp_unref_tgt ((void *) k
->tgt
);
1384 return is_tgt_unmapped
;
1387 attribute_hidden
bool
1388 gomp_remove_var (struct gomp_device_descr
*devicep
, splay_tree_key k
)
1390 return gomp_remove_var_internal (devicep
, k
, NULL
);
1393 /* Remove a variable asynchronously. This actually removes the variable
1394 mapping immediately, but retains the linked target_mem_desc until the
1395 asynchronous operation has completed (as it may still refer to target
1396 memory). The device lock must be held before entry, and remains locked on
1399 attribute_hidden
void
1400 gomp_remove_var_async (struct gomp_device_descr
*devicep
, splay_tree_key k
,
1401 struct goacc_asyncqueue
*aq
)
1403 (void) gomp_remove_var_internal (devicep
, k
, aq
);
1406 /* Unmap variables described by TGT. If DO_COPYFROM is true, copy relevant
1407 variables back from device to host: if it is false, it is assumed that this
1408 has been done already. */
1410 static inline __attribute__((always_inline
)) void
1411 gomp_unmap_vars_internal (struct target_mem_desc
*tgt
, bool do_copyfrom
,
1412 struct goacc_asyncqueue
*aq
)
1414 struct gomp_device_descr
*devicep
= tgt
->device_descr
;
1416 if (tgt
->list_count
== 0)
1422 gomp_mutex_lock (&devicep
->lock
);
1423 if (devicep
->state
== GOMP_DEVICE_FINALIZED
)
1425 gomp_mutex_unlock (&devicep
->lock
);
1433 /* We must perform detachments before any copies back to the host. */
1434 for (i
= 0; i
< tgt
->list_count
; i
++)
1436 splay_tree_key k
= tgt
->list
[i
].key
;
1438 if (k
!= NULL
&& tgt
->list
[i
].is_attach
)
1439 gomp_detach_pointer (devicep
, aq
, k
, tgt
->list
[i
].key
->host_start
1440 + tgt
->list
[i
].offset
,
1444 for (i
= 0; i
< tgt
->list_count
; i
++)
1446 splay_tree_key k
= tgt
->list
[i
].key
;
1450 /* OpenACC 'attach'/'detach' doesn't affect structured/dynamic reference
1451 counts ('n->refcount', 'n->dynamic_refcount'). */
1452 if (tgt
->list
[i
].is_attach
)
1455 bool do_unmap
= false;
1456 if (k
->refcount
> 1 && k
->refcount
!= REFCOUNT_INFINITY
)
1458 else if (k
->refcount
== 1)
1464 if ((do_unmap
&& do_copyfrom
&& tgt
->list
[i
].copy_from
)
1465 || tgt
->list
[i
].always_copy_from
)
1466 gomp_copy_dev2host (devicep
, aq
,
1467 (void *) (k
->host_start
+ tgt
->list
[i
].offset
),
1468 (void *) (k
->tgt
->tgt_start
+ k
->tgt_offset
1469 + tgt
->list
[i
].offset
),
1470 tgt
->list
[i
].length
);
1473 struct target_mem_desc
*k_tgt
= k
->tgt
;
1474 bool is_tgt_unmapped
= gomp_remove_var (devicep
, k
);
1475 /* It would be bad if TGT got unmapped while we're still iterating
1476 over its LIST_COUNT, and also expect to use it in the following
1478 assert (!is_tgt_unmapped
1484 devicep
->openacc
.async
.queue_callback_func (aq
, gomp_unref_tgt_void
,
1487 gomp_unref_tgt ((void *) tgt
);
1489 gomp_mutex_unlock (&devicep
->lock
);
1492 attribute_hidden
void
1493 gomp_unmap_vars (struct target_mem_desc
*tgt
, bool do_copyfrom
)
1495 gomp_unmap_vars_internal (tgt
, do_copyfrom
, NULL
);
1498 attribute_hidden
void
1499 gomp_unmap_vars_async (struct target_mem_desc
*tgt
, bool do_copyfrom
,
1500 struct goacc_asyncqueue
*aq
)
1502 gomp_unmap_vars_internal (tgt
, do_copyfrom
, aq
);
1506 gomp_update (struct gomp_device_descr
*devicep
, size_t mapnum
, void **hostaddrs
,
1507 size_t *sizes
, void *kinds
, bool short_mapkind
)
1510 struct splay_tree_key_s cur_node
;
1511 const int typemask
= short_mapkind
? 0xff : 0x7;
1519 gomp_mutex_lock (&devicep
->lock
);
1520 if (devicep
->state
== GOMP_DEVICE_FINALIZED
)
1522 gomp_mutex_unlock (&devicep
->lock
);
1526 for (i
= 0; i
< mapnum
; i
++)
1529 cur_node
.host_start
= (uintptr_t) hostaddrs
[i
];
1530 cur_node
.host_end
= cur_node
.host_start
+ sizes
[i
];
1531 splay_tree_key n
= splay_tree_lookup (&devicep
->mem_map
, &cur_node
);
1534 int kind
= get_kind (short_mapkind
, kinds
, i
);
1535 if (n
->host_start
> cur_node
.host_start
1536 || n
->host_end
< cur_node
.host_end
)
1538 gomp_mutex_unlock (&devicep
->lock
);
1539 gomp_fatal ("Trying to update [%p..%p) object when "
1540 "only [%p..%p) is mapped",
1541 (void *) cur_node
.host_start
,
1542 (void *) cur_node
.host_end
,
1543 (void *) n
->host_start
,
1544 (void *) n
->host_end
);
1548 void *hostaddr
= (void *) cur_node
.host_start
;
1549 void *devaddr
= (void *) (n
->tgt
->tgt_start
+ n
->tgt_offset
1550 + cur_node
.host_start
- n
->host_start
);
1551 size_t size
= cur_node
.host_end
- cur_node
.host_start
;
1553 if (GOMP_MAP_COPY_TO_P (kind
& typemask
))
1554 gomp_copy_host2dev (devicep
, NULL
, devaddr
, hostaddr
, size
,
1556 if (GOMP_MAP_COPY_FROM_P (kind
& typemask
))
1557 gomp_copy_dev2host (devicep
, NULL
, hostaddr
, devaddr
, size
);
1560 gomp_mutex_unlock (&devicep
->lock
);
1563 /* Load image pointed by TARGET_DATA to the device, specified by DEVICEP.
1564 And insert to splay tree the mapping between addresses from HOST_TABLE and
1565 from loaded target image. We rely in the host and device compiler
1566 emitting variable and functions in the same order. */
1569 gomp_load_image_to_device (struct gomp_device_descr
*devicep
, unsigned version
,
1570 const void *host_table
, const void *target_data
,
1571 bool is_register_lock
)
1573 void **host_func_table
= ((void ***) host_table
)[0];
1574 void **host_funcs_end
= ((void ***) host_table
)[1];
1575 void **host_var_table
= ((void ***) host_table
)[2];
1576 void **host_vars_end
= ((void ***) host_table
)[3];
1578 /* The func table contains only addresses, the var table contains addresses
1579 and corresponding sizes. */
1580 int num_funcs
= host_funcs_end
- host_func_table
;
1581 int num_vars
= (host_vars_end
- host_var_table
) / 2;
1583 /* Load image to device and get target addresses for the image. */
1584 struct addr_pair
*target_table
= NULL
;
1585 int i
, num_target_entries
;
1588 = devicep
->load_image_func (devicep
->target_id
, version
,
1589 target_data
, &target_table
);
1591 if (num_target_entries
!= num_funcs
+ num_vars
)
1593 gomp_mutex_unlock (&devicep
->lock
);
1594 if (is_register_lock
)
1595 gomp_mutex_unlock (®ister_lock
);
1596 gomp_fatal ("Cannot map target functions or variables"
1597 " (expected %u, have %u)", num_funcs
+ num_vars
,
1598 num_target_entries
);
1601 /* Insert host-target address mapping into splay tree. */
1602 struct target_mem_desc
*tgt
= gomp_malloc (sizeof (*tgt
));
1603 tgt
->array
= gomp_malloc ((num_funcs
+ num_vars
) * sizeof (*tgt
->array
));
1604 tgt
->refcount
= REFCOUNT_INFINITY
;
1607 tgt
->to_free
= NULL
;
1609 tgt
->list_count
= 0;
1610 tgt
->device_descr
= devicep
;
1611 splay_tree_node array
= tgt
->array
;
1613 for (i
= 0; i
< num_funcs
; i
++)
1615 splay_tree_key k
= &array
->key
;
1616 k
->host_start
= (uintptr_t) host_func_table
[i
];
1617 k
->host_end
= k
->host_start
+ 1;
1619 k
->tgt_offset
= target_table
[i
].start
;
1620 k
->refcount
= REFCOUNT_INFINITY
;
1621 k
->dynamic_refcount
= 0;
1624 array
->right
= NULL
;
1625 splay_tree_insert (&devicep
->mem_map
, array
);
1629 /* Most significant bit of the size in host and target tables marks
1630 "omp declare target link" variables. */
1631 const uintptr_t link_bit
= 1ULL << (sizeof (uintptr_t) * __CHAR_BIT__
- 1);
1632 const uintptr_t size_mask
= ~link_bit
;
1634 for (i
= 0; i
< num_vars
; i
++)
1636 struct addr_pair
*target_var
= &target_table
[num_funcs
+ i
];
1637 uintptr_t target_size
= target_var
->end
- target_var
->start
;
1638 bool is_link_var
= link_bit
& (uintptr_t) host_var_table
[i
* 2 + 1];
1640 if (!is_link_var
&& (uintptr_t) host_var_table
[i
* 2 + 1] != target_size
)
1642 gomp_mutex_unlock (&devicep
->lock
);
1643 if (is_register_lock
)
1644 gomp_mutex_unlock (®ister_lock
);
1645 gomp_fatal ("Cannot map target variables (size mismatch)");
1648 splay_tree_key k
= &array
->key
;
1649 k
->host_start
= (uintptr_t) host_var_table
[i
* 2];
1651 = k
->host_start
+ (size_mask
& (uintptr_t) host_var_table
[i
* 2 + 1]);
1653 k
->tgt_offset
= target_var
->start
;
1654 k
->refcount
= is_link_var
? REFCOUNT_LINK
: REFCOUNT_INFINITY
;
1655 k
->dynamic_refcount
= 0;
1658 array
->right
= NULL
;
1659 splay_tree_insert (&devicep
->mem_map
, array
);
1663 free (target_table
);
1666 /* Unload the mappings described by target_data from device DEVICE_P.
1667 The device must be locked. */
1670 gomp_unload_image_from_device (struct gomp_device_descr
*devicep
,
1672 const void *host_table
, const void *target_data
)
1674 void **host_func_table
= ((void ***) host_table
)[0];
1675 void **host_funcs_end
= ((void ***) host_table
)[1];
1676 void **host_var_table
= ((void ***) host_table
)[2];
1677 void **host_vars_end
= ((void ***) host_table
)[3];
1679 /* The func table contains only addresses, the var table contains addresses
1680 and corresponding sizes. */
1681 int num_funcs
= host_funcs_end
- host_func_table
;
1682 int num_vars
= (host_vars_end
- host_var_table
) / 2;
1684 struct splay_tree_key_s k
;
1685 splay_tree_key node
= NULL
;
1687 /* Find mapping at start of node array */
1688 if (num_funcs
|| num_vars
)
1690 k
.host_start
= (num_funcs
? (uintptr_t) host_func_table
[0]
1691 : (uintptr_t) host_var_table
[0]);
1692 k
.host_end
= k
.host_start
+ 1;
1693 node
= splay_tree_lookup (&devicep
->mem_map
, &k
);
1696 if (!devicep
->unload_image_func (devicep
->target_id
, version
, target_data
))
1698 gomp_mutex_unlock (&devicep
->lock
);
1699 gomp_fatal ("image unload fail");
1702 /* Remove mappings from splay tree. */
1704 for (i
= 0; i
< num_funcs
; i
++)
1706 k
.host_start
= (uintptr_t) host_func_table
[i
];
1707 k
.host_end
= k
.host_start
+ 1;
1708 splay_tree_remove (&devicep
->mem_map
, &k
);
1711 /* Most significant bit of the size in host and target tables marks
1712 "omp declare target link" variables. */
1713 const uintptr_t link_bit
= 1ULL << (sizeof (uintptr_t) * __CHAR_BIT__
- 1);
1714 const uintptr_t size_mask
= ~link_bit
;
1715 bool is_tgt_unmapped
= false;
1717 for (i
= 0; i
< num_vars
; i
++)
1719 k
.host_start
= (uintptr_t) host_var_table
[i
* 2];
1721 = k
.host_start
+ (size_mask
& (uintptr_t) host_var_table
[i
* 2 + 1]);
1723 if (!(link_bit
& (uintptr_t) host_var_table
[i
* 2 + 1]))
1724 splay_tree_remove (&devicep
->mem_map
, &k
);
1727 splay_tree_key n
= splay_tree_lookup (&devicep
->mem_map
, &k
);
1728 is_tgt_unmapped
= gomp_remove_var (devicep
, n
);
1732 if (node
&& !is_tgt_unmapped
)
1739 /* This function should be called from every offload image while loading.
1740 It gets the descriptor of the host func and var tables HOST_TABLE, TYPE of
1741 the target, and TARGET_DATA needed by target plugin. */
1744 GOMP_offload_register_ver (unsigned version
, const void *host_table
,
1745 int target_type
, const void *target_data
)
1749 if (GOMP_VERSION_LIB (version
) > GOMP_VERSION
)
1750 gomp_fatal ("Library too old for offload (version %u < %u)",
1751 GOMP_VERSION
, GOMP_VERSION_LIB (version
));
1753 gomp_mutex_lock (®ister_lock
);
1755 /* Load image to all initialized devices. */
1756 for (i
= 0; i
< num_devices
; i
++)
1758 struct gomp_device_descr
*devicep
= &devices
[i
];
1759 gomp_mutex_lock (&devicep
->lock
);
1760 if (devicep
->type
== target_type
1761 && devicep
->state
== GOMP_DEVICE_INITIALIZED
)
1762 gomp_load_image_to_device (devicep
, version
,
1763 host_table
, target_data
, true);
1764 gomp_mutex_unlock (&devicep
->lock
);
1767 /* Insert image to array of pending images. */
1769 = gomp_realloc_unlock (offload_images
,
1770 (num_offload_images
+ 1)
1771 * sizeof (struct offload_image_descr
));
1772 offload_images
[num_offload_images
].version
= version
;
1773 offload_images
[num_offload_images
].type
= target_type
;
1774 offload_images
[num_offload_images
].host_table
= host_table
;
1775 offload_images
[num_offload_images
].target_data
= target_data
;
1777 num_offload_images
++;
1778 gomp_mutex_unlock (®ister_lock
);
1782 GOMP_offload_register (const void *host_table
, int target_type
,
1783 const void *target_data
)
1785 GOMP_offload_register_ver (0, host_table
, target_type
, target_data
);
1788 /* This function should be called from every offload image while unloading.
1789 It gets the descriptor of the host func and var tables HOST_TABLE, TYPE of
1790 the target, and TARGET_DATA needed by target plugin. */
1793 GOMP_offload_unregister_ver (unsigned version
, const void *host_table
,
1794 int target_type
, const void *target_data
)
1798 gomp_mutex_lock (®ister_lock
);
1800 /* Unload image from all initialized devices. */
1801 for (i
= 0; i
< num_devices
; i
++)
1803 struct gomp_device_descr
*devicep
= &devices
[i
];
1804 gomp_mutex_lock (&devicep
->lock
);
1805 if (devicep
->type
== target_type
1806 && devicep
->state
== GOMP_DEVICE_INITIALIZED
)
1807 gomp_unload_image_from_device (devicep
, version
,
1808 host_table
, target_data
);
1809 gomp_mutex_unlock (&devicep
->lock
);
1812 /* Remove image from array of pending images. */
1813 for (i
= 0; i
< num_offload_images
; i
++)
1814 if (offload_images
[i
].target_data
== target_data
)
1816 offload_images
[i
] = offload_images
[--num_offload_images
];
1820 gomp_mutex_unlock (®ister_lock
);
1824 GOMP_offload_unregister (const void *host_table
, int target_type
,
1825 const void *target_data
)
1827 GOMP_offload_unregister_ver (0, host_table
, target_type
, target_data
);
1830 /* This function initializes the target device, specified by DEVICEP. DEVICEP
1831 must be locked on entry, and remains locked on return. */
1833 attribute_hidden
void
1834 gomp_init_device (struct gomp_device_descr
*devicep
)
1837 if (!devicep
->init_device_func (devicep
->target_id
))
1839 gomp_mutex_unlock (&devicep
->lock
);
1840 gomp_fatal ("device initialization failed");
1843 /* Load to device all images registered by the moment. */
1844 for (i
= 0; i
< num_offload_images
; i
++)
1846 struct offload_image_descr
*image
= &offload_images
[i
];
1847 if (image
->type
== devicep
->type
)
1848 gomp_load_image_to_device (devicep
, image
->version
,
1849 image
->host_table
, image
->target_data
,
1853 /* Initialize OpenACC asynchronous queues. */
1854 goacc_init_asyncqueues (devicep
);
1856 devicep
->state
= GOMP_DEVICE_INITIALIZED
;
1859 /* This function finalizes the target device, specified by DEVICEP. DEVICEP
1860 must be locked on entry, and remains locked on return. */
1862 attribute_hidden
bool
1863 gomp_fini_device (struct gomp_device_descr
*devicep
)
1865 bool ret
= goacc_fini_asyncqueues (devicep
);
1866 ret
&= devicep
->fini_device_func (devicep
->target_id
);
1867 devicep
->state
= GOMP_DEVICE_FINALIZED
;
1871 attribute_hidden
void
1872 gomp_unload_device (struct gomp_device_descr
*devicep
)
1874 if (devicep
->state
== GOMP_DEVICE_INITIALIZED
)
1878 /* Unload from device all images registered at the moment. */
1879 for (i
= 0; i
< num_offload_images
; i
++)
1881 struct offload_image_descr
*image
= &offload_images
[i
];
1882 if (image
->type
== devicep
->type
)
1883 gomp_unload_image_from_device (devicep
, image
->version
,
1885 image
->target_data
);
1890 /* Host fallback for GOMP_target{,_ext} routines. */
1893 gomp_target_fallback (void (*fn
) (void *), void **hostaddrs
)
1895 struct gomp_thread old_thr
, *thr
= gomp_thread ();
1897 memset (thr
, '\0', sizeof (*thr
));
1898 if (gomp_places_list
)
1900 thr
->place
= old_thr
.place
;
1901 thr
->ts
.place_partition_len
= gomp_places_list_len
;
1904 gomp_free_thread (thr
);
1908 /* Calculate alignment and size requirements of a private copy of data shared
1909 as GOMP_MAP_FIRSTPRIVATE and store them to TGT_ALIGN and TGT_SIZE. */
1912 calculate_firstprivate_requirements (size_t mapnum
, size_t *sizes
,
1913 unsigned short *kinds
, size_t *tgt_align
,
1917 for (i
= 0; i
< mapnum
; i
++)
1918 if ((kinds
[i
] & 0xff) == GOMP_MAP_FIRSTPRIVATE
)
1920 size_t align
= (size_t) 1 << (kinds
[i
] >> 8);
1921 if (*tgt_align
< align
)
1923 *tgt_size
= (*tgt_size
+ align
- 1) & ~(align
- 1);
1924 *tgt_size
+= sizes
[i
];
1928 /* Copy data shared as GOMP_MAP_FIRSTPRIVATE to DST. */
1931 copy_firstprivate_data (char *tgt
, size_t mapnum
, void **hostaddrs
,
1932 size_t *sizes
, unsigned short *kinds
, size_t tgt_align
,
1935 uintptr_t al
= (uintptr_t) tgt
& (tgt_align
- 1);
1937 tgt
+= tgt_align
- al
;
1940 for (i
= 0; i
< mapnum
; i
++)
1941 if ((kinds
[i
] & 0xff) == GOMP_MAP_FIRSTPRIVATE
)
1943 size_t align
= (size_t) 1 << (kinds
[i
] >> 8);
1944 tgt_size
= (tgt_size
+ align
- 1) & ~(align
- 1);
1945 memcpy (tgt
+ tgt_size
, hostaddrs
[i
], sizes
[i
]);
1946 hostaddrs
[i
] = tgt
+ tgt_size
;
1947 tgt_size
= tgt_size
+ sizes
[i
];
1951 /* Helper function of GOMP_target{,_ext} routines. */
1954 gomp_get_target_fn_addr (struct gomp_device_descr
*devicep
,
1955 void (*host_fn
) (void *))
1957 if (devicep
->capabilities
& GOMP_OFFLOAD_CAP_NATIVE_EXEC
)
1958 return (void *) host_fn
;
1961 gomp_mutex_lock (&devicep
->lock
);
1962 if (devicep
->state
== GOMP_DEVICE_FINALIZED
)
1964 gomp_mutex_unlock (&devicep
->lock
);
1968 struct splay_tree_key_s k
;
1969 k
.host_start
= (uintptr_t) host_fn
;
1970 k
.host_end
= k
.host_start
+ 1;
1971 splay_tree_key tgt_fn
= splay_tree_lookup (&devicep
->mem_map
, &k
);
1972 gomp_mutex_unlock (&devicep
->lock
);
1976 return (void *) tgt_fn
->tgt_offset
;
1980 /* Called when encountering a target directive. If DEVICE
1981 is GOMP_DEVICE_ICV, it means use device-var ICV. If it is
1982 GOMP_DEVICE_HOST_FALLBACK (or any value
1983 larger than last available hw device), use host fallback.
1984 FN is address of host code, UNUSED is part of the current ABI, but
1985 we're not actually using it. HOSTADDRS, SIZES and KINDS are arrays
1986 with MAPNUM entries, with addresses of the host objects,
1987 sizes of the host objects (resp. for pointer kind pointer bias
1988 and assumed sizeof (void *) size) and kinds. */
1991 GOMP_target (int device
, void (*fn
) (void *), const void *unused
,
1992 size_t mapnum
, void **hostaddrs
, size_t *sizes
,
1993 unsigned char *kinds
)
1995 struct gomp_device_descr
*devicep
= resolve_device (device
);
1999 || !(devicep
->capabilities
& GOMP_OFFLOAD_CAP_OPENMP_400
)
2000 /* All shared memory devices should use the GOMP_target_ext function. */
2001 || devicep
->capabilities
& GOMP_OFFLOAD_CAP_SHARED_MEM
2002 || !(fn_addr
= gomp_get_target_fn_addr (devicep
, fn
)))
2003 return gomp_target_fallback (fn
, hostaddrs
);
2005 struct target_mem_desc
*tgt_vars
2006 = gomp_map_vars (devicep
, mapnum
, hostaddrs
, NULL
, sizes
, kinds
, false,
2007 GOMP_MAP_VARS_TARGET
);
2008 devicep
->run_func (devicep
->target_id
, fn_addr
, (void *) tgt_vars
->tgt_start
,
2010 gomp_unmap_vars (tgt_vars
, true);
2013 static inline unsigned int
2014 clear_unsupported_flags (struct gomp_device_descr
*devicep
, unsigned int flags
)
2016 /* If we cannot run asynchronously, simply ignore nowait. */
2017 if (devicep
!= NULL
&& devicep
->async_run_func
== NULL
)
2018 flags
&= ~GOMP_TARGET_FLAG_NOWAIT
;
2023 /* Like GOMP_target, but KINDS is 16-bit, UNUSED is no longer present,
2024 and several arguments have been added:
2025 FLAGS is a bitmask, see GOMP_TARGET_FLAG_* in gomp-constants.h.
2026 DEPEND is array of dependencies, see GOMP_task for details.
2028 ARGS is a pointer to an array consisting of a variable number of both
2029 device-independent and device-specific arguments, which can take one two
2030 elements where the first specifies for which device it is intended, the type
2031 and optionally also the value. If the value is not present in the first
2032 one, the whole second element the actual value. The last element of the
2033 array is a single NULL. Among the device independent can be for example
2034 NUM_TEAMS and THREAD_LIMIT.
2036 NUM_TEAMS is positive if GOMP_teams will be called in the body with
2037 that value, or 1 if teams construct is not present, or 0, if
2038 teams construct does not have num_teams clause and so the choice is
2039 implementation defined, and -1 if it can't be determined on the host
2040 what value will GOMP_teams have on the device.
2041 THREAD_LIMIT similarly is positive if GOMP_teams will be called in the
2042 body with that value, or 0, if teams construct does not have thread_limit
2043 clause or the teams construct is not present, or -1 if it can't be
2044 determined on the host what value will GOMP_teams have on the device. */
2047 GOMP_target_ext (int device
, void (*fn
) (void *), size_t mapnum
,
2048 void **hostaddrs
, size_t *sizes
, unsigned short *kinds
,
2049 unsigned int flags
, void **depend
, void **args
)
2051 struct gomp_device_descr
*devicep
= resolve_device (device
);
2052 size_t tgt_align
= 0, tgt_size
= 0;
2053 bool fpc_done
= false;
2055 flags
= clear_unsupported_flags (devicep
, flags
);
2057 if (flags
& GOMP_TARGET_FLAG_NOWAIT
)
2059 struct gomp_thread
*thr
= gomp_thread ();
2060 /* Create a team if we don't have any around, as nowait
2061 target tasks make sense to run asynchronously even when
2062 outside of any parallel. */
2063 if (__builtin_expect (thr
->ts
.team
== NULL
, 0))
2065 struct gomp_team
*team
= gomp_new_team (1);
2066 struct gomp_task
*task
= thr
->task
;
2067 struct gomp_task_icv
*icv
= task
? &task
->icv
: &gomp_global_icv
;
2068 team
->prev_ts
= thr
->ts
;
2069 thr
->ts
.team
= team
;
2070 thr
->ts
.team_id
= 0;
2071 thr
->ts
.work_share
= &team
->work_shares
[0];
2072 thr
->ts
.last_work_share
= NULL
;
2073 #ifdef HAVE_SYNC_BUILTINS
2074 thr
->ts
.single_count
= 0;
2076 thr
->ts
.static_trip
= 0;
2077 thr
->task
= &team
->implicit_task
[0];
2078 gomp_init_task (thr
->task
, NULL
, icv
);
2084 thr
->task
= &team
->implicit_task
[0];
2087 pthread_setspecific (gomp_thread_destructor
, thr
);
2090 && !thr
->task
->final_task
)
2092 gomp_create_target_task (devicep
, fn
, mapnum
, hostaddrs
,
2093 sizes
, kinds
, flags
, depend
, args
,
2094 GOMP_TARGET_TASK_BEFORE_MAP
);
2099 /* If there are depend clauses, but nowait is not present
2100 (or we are in a final task), block the parent task until the
2101 dependencies are resolved and then just continue with the rest
2102 of the function as if it is a merged task. */
2105 struct gomp_thread
*thr
= gomp_thread ();
2106 if (thr
->task
&& thr
->task
->depend_hash
)
2108 /* If we might need to wait, copy firstprivate now. */
2109 calculate_firstprivate_requirements (mapnum
, sizes
, kinds
,
2110 &tgt_align
, &tgt_size
);
2113 char *tgt
= gomp_alloca (tgt_size
+ tgt_align
- 1);
2114 copy_firstprivate_data (tgt
, mapnum
, hostaddrs
, sizes
, kinds
,
2115 tgt_align
, tgt_size
);
2118 gomp_task_maybe_wait_for_dependencies (depend
);
2124 || !(devicep
->capabilities
& GOMP_OFFLOAD_CAP_OPENMP_400
)
2125 || !(fn_addr
= gomp_get_target_fn_addr (devicep
, fn
))
2126 || (devicep
->can_run_func
&& !devicep
->can_run_func (fn_addr
)))
2130 calculate_firstprivate_requirements (mapnum
, sizes
, kinds
,
2131 &tgt_align
, &tgt_size
);
2134 char *tgt
= gomp_alloca (tgt_size
+ tgt_align
- 1);
2135 copy_firstprivate_data (tgt
, mapnum
, hostaddrs
, sizes
, kinds
,
2136 tgt_align
, tgt_size
);
2139 gomp_target_fallback (fn
, hostaddrs
);
2143 struct target_mem_desc
*tgt_vars
;
2144 if (devicep
->capabilities
& GOMP_OFFLOAD_CAP_SHARED_MEM
)
2148 calculate_firstprivate_requirements (mapnum
, sizes
, kinds
,
2149 &tgt_align
, &tgt_size
);
2152 char *tgt
= gomp_alloca (tgt_size
+ tgt_align
- 1);
2153 copy_firstprivate_data (tgt
, mapnum
, hostaddrs
, sizes
, kinds
,
2154 tgt_align
, tgt_size
);
2160 tgt_vars
= gomp_map_vars (devicep
, mapnum
, hostaddrs
, NULL
, sizes
, kinds
,
2161 true, GOMP_MAP_VARS_TARGET
);
2162 devicep
->run_func (devicep
->target_id
, fn_addr
,
2163 tgt_vars
? (void *) tgt_vars
->tgt_start
: hostaddrs
,
2166 gomp_unmap_vars (tgt_vars
, true);
2169 /* Host fallback for GOMP_target_data{,_ext} routines. */
2172 gomp_target_data_fallback (void)
2174 struct gomp_task_icv
*icv
= gomp_icv (false);
2175 if (icv
->target_data
)
2177 /* Even when doing a host fallback, if there are any active
2178 #pragma omp target data constructs, need to remember the
2179 new #pragma omp target data, otherwise GOMP_target_end_data
2180 would get out of sync. */
2181 struct target_mem_desc
*tgt
2182 = gomp_map_vars (NULL
, 0, NULL
, NULL
, NULL
, NULL
, false,
2183 GOMP_MAP_VARS_DATA
);
2184 tgt
->prev
= icv
->target_data
;
2185 icv
->target_data
= tgt
;
2190 GOMP_target_data (int device
, const void *unused
, size_t mapnum
,
2191 void **hostaddrs
, size_t *sizes
, unsigned char *kinds
)
2193 struct gomp_device_descr
*devicep
= resolve_device (device
);
2196 || !(devicep
->capabilities
& GOMP_OFFLOAD_CAP_OPENMP_400
)
2197 || (devicep
->capabilities
& GOMP_OFFLOAD_CAP_SHARED_MEM
))
2198 return gomp_target_data_fallback ();
2200 struct target_mem_desc
*tgt
2201 = gomp_map_vars (devicep
, mapnum
, hostaddrs
, NULL
, sizes
, kinds
, false,
2202 GOMP_MAP_VARS_DATA
);
2203 struct gomp_task_icv
*icv
= gomp_icv (true);
2204 tgt
->prev
= icv
->target_data
;
2205 icv
->target_data
= tgt
;
2209 GOMP_target_data_ext (int device
, size_t mapnum
, void **hostaddrs
,
2210 size_t *sizes
, unsigned short *kinds
)
2212 struct gomp_device_descr
*devicep
= resolve_device (device
);
2215 || !(devicep
->capabilities
& GOMP_OFFLOAD_CAP_OPENMP_400
)
2216 || devicep
->capabilities
& GOMP_OFFLOAD_CAP_SHARED_MEM
)
2217 return gomp_target_data_fallback ();
2219 struct target_mem_desc
*tgt
2220 = gomp_map_vars (devicep
, mapnum
, hostaddrs
, NULL
, sizes
, kinds
, true,
2221 GOMP_MAP_VARS_DATA
);
2222 struct gomp_task_icv
*icv
= gomp_icv (true);
2223 tgt
->prev
= icv
->target_data
;
2224 icv
->target_data
= tgt
;
2228 GOMP_target_end_data (void)
2230 struct gomp_task_icv
*icv
= gomp_icv (false);
2231 if (icv
->target_data
)
2233 struct target_mem_desc
*tgt
= icv
->target_data
;
2234 icv
->target_data
= tgt
->prev
;
2235 gomp_unmap_vars (tgt
, true);
2240 GOMP_target_update (int device
, const void *unused
, size_t mapnum
,
2241 void **hostaddrs
, size_t *sizes
, unsigned char *kinds
)
2243 struct gomp_device_descr
*devicep
= resolve_device (device
);
2246 || !(devicep
->capabilities
& GOMP_OFFLOAD_CAP_OPENMP_400
)
2247 || devicep
->capabilities
& GOMP_OFFLOAD_CAP_SHARED_MEM
)
2250 gomp_update (devicep
, mapnum
, hostaddrs
, sizes
, kinds
, false);
2254 GOMP_target_update_ext (int device
, size_t mapnum
, void **hostaddrs
,
2255 size_t *sizes
, unsigned short *kinds
,
2256 unsigned int flags
, void **depend
)
2258 struct gomp_device_descr
*devicep
= resolve_device (device
);
2260 /* If there are depend clauses, but nowait is not present,
2261 block the parent task until the dependencies are resolved
2262 and then just continue with the rest of the function as if it
2263 is a merged task. Until we are able to schedule task during
2264 variable mapping or unmapping, ignore nowait if depend clauses
2268 struct gomp_thread
*thr
= gomp_thread ();
2269 if (thr
->task
&& thr
->task
->depend_hash
)
2271 if ((flags
& GOMP_TARGET_FLAG_NOWAIT
)
2273 && !thr
->task
->final_task
)
2275 if (gomp_create_target_task (devicep
, (void (*) (void *)) NULL
,
2276 mapnum
, hostaddrs
, sizes
, kinds
,
2277 flags
| GOMP_TARGET_FLAG_UPDATE
,
2278 depend
, NULL
, GOMP_TARGET_TASK_DATA
))
2283 struct gomp_team
*team
= thr
->ts
.team
;
2284 /* If parallel or taskgroup has been cancelled, don't start new
2286 if (__builtin_expect (gomp_cancel_var
, 0) && team
)
2288 if (gomp_team_barrier_cancelled (&team
->barrier
))
2290 if (thr
->task
->taskgroup
)
2292 if (thr
->task
->taskgroup
->cancelled
)
2294 if (thr
->task
->taskgroup
->workshare
2295 && thr
->task
->taskgroup
->prev
2296 && thr
->task
->taskgroup
->prev
->cancelled
)
2301 gomp_task_maybe_wait_for_dependencies (depend
);
2307 || !(devicep
->capabilities
& GOMP_OFFLOAD_CAP_OPENMP_400
)
2308 || devicep
->capabilities
& GOMP_OFFLOAD_CAP_SHARED_MEM
)
2311 struct gomp_thread
*thr
= gomp_thread ();
2312 struct gomp_team
*team
= thr
->ts
.team
;
2313 /* If parallel or taskgroup has been cancelled, don't start new tasks. */
2314 if (__builtin_expect (gomp_cancel_var
, 0) && team
)
2316 if (gomp_team_barrier_cancelled (&team
->barrier
))
2318 if (thr
->task
->taskgroup
)
2320 if (thr
->task
->taskgroup
->cancelled
)
2322 if (thr
->task
->taskgroup
->workshare
2323 && thr
->task
->taskgroup
->prev
2324 && thr
->task
->taskgroup
->prev
->cancelled
)
2329 gomp_update (devicep
, mapnum
, hostaddrs
, sizes
, kinds
, true);
2333 gomp_exit_data (struct gomp_device_descr
*devicep
, size_t mapnum
,
2334 void **hostaddrs
, size_t *sizes
, unsigned short *kinds
)
2336 const int typemask
= 0xff;
2338 gomp_mutex_lock (&devicep
->lock
);
2339 if (devicep
->state
== GOMP_DEVICE_FINALIZED
)
2341 gomp_mutex_unlock (&devicep
->lock
);
2345 for (i
= 0; i
< mapnum
; i
++)
2347 struct splay_tree_key_s cur_node
;
2348 unsigned char kind
= kinds
[i
] & typemask
;
2352 case GOMP_MAP_ALWAYS_FROM
:
2353 case GOMP_MAP_DELETE
:
2354 case GOMP_MAP_RELEASE
:
2355 case GOMP_MAP_ZERO_LEN_ARRAY_SECTION
:
2356 case GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION
:
2357 cur_node
.host_start
= (uintptr_t) hostaddrs
[i
];
2358 cur_node
.host_end
= cur_node
.host_start
+ sizes
[i
];
2359 splay_tree_key k
= (kind
== GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION
2360 || kind
== GOMP_MAP_ZERO_LEN_ARRAY_SECTION
)
2361 ? gomp_map_0len_lookup (&devicep
->mem_map
, &cur_node
)
2362 : splay_tree_lookup (&devicep
->mem_map
, &cur_node
);
2366 if (k
->refcount
> 0 && k
->refcount
!= REFCOUNT_INFINITY
)
2368 if ((kind
== GOMP_MAP_DELETE
2369 || kind
== GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION
)
2370 && k
->refcount
!= REFCOUNT_INFINITY
)
2373 if ((kind
== GOMP_MAP_FROM
&& k
->refcount
== 0)
2374 || kind
== GOMP_MAP_ALWAYS_FROM
)
2375 gomp_copy_dev2host (devicep
, NULL
, (void *) cur_node
.host_start
,
2376 (void *) (k
->tgt
->tgt_start
+ k
->tgt_offset
2377 + cur_node
.host_start
2379 cur_node
.host_end
- cur_node
.host_start
);
2380 if (k
->refcount
== 0)
2381 gomp_remove_var (devicep
, k
);
2385 gomp_mutex_unlock (&devicep
->lock
);
2386 gomp_fatal ("GOMP_target_enter_exit_data unhandled kind 0x%.2x",
2391 gomp_mutex_unlock (&devicep
->lock
);
2395 GOMP_target_enter_exit_data (int device
, size_t mapnum
, void **hostaddrs
,
2396 size_t *sizes
, unsigned short *kinds
,
2397 unsigned int flags
, void **depend
)
2399 struct gomp_device_descr
*devicep
= resolve_device (device
);
2401 /* If there are depend clauses, but nowait is not present,
2402 block the parent task until the dependencies are resolved
2403 and then just continue with the rest of the function as if it
2404 is a merged task. Until we are able to schedule task during
2405 variable mapping or unmapping, ignore nowait if depend clauses
2409 struct gomp_thread
*thr
= gomp_thread ();
2410 if (thr
->task
&& thr
->task
->depend_hash
)
2412 if ((flags
& GOMP_TARGET_FLAG_NOWAIT
)
2414 && !thr
->task
->final_task
)
2416 if (gomp_create_target_task (devicep
, (void (*) (void *)) NULL
,
2417 mapnum
, hostaddrs
, sizes
, kinds
,
2418 flags
, depend
, NULL
,
2419 GOMP_TARGET_TASK_DATA
))
2424 struct gomp_team
*team
= thr
->ts
.team
;
2425 /* If parallel or taskgroup has been cancelled, don't start new
2427 if (__builtin_expect (gomp_cancel_var
, 0) && team
)
2429 if (gomp_team_barrier_cancelled (&team
->barrier
))
2431 if (thr
->task
->taskgroup
)
2433 if (thr
->task
->taskgroup
->cancelled
)
2435 if (thr
->task
->taskgroup
->workshare
2436 && thr
->task
->taskgroup
->prev
2437 && thr
->task
->taskgroup
->prev
->cancelled
)
2442 gomp_task_maybe_wait_for_dependencies (depend
);
2448 || !(devicep
->capabilities
& GOMP_OFFLOAD_CAP_OPENMP_400
)
2449 || devicep
->capabilities
& GOMP_OFFLOAD_CAP_SHARED_MEM
)
2452 struct gomp_thread
*thr
= gomp_thread ();
2453 struct gomp_team
*team
= thr
->ts
.team
;
2454 /* If parallel or taskgroup has been cancelled, don't start new tasks. */
2455 if (__builtin_expect (gomp_cancel_var
, 0) && team
)
2457 if (gomp_team_barrier_cancelled (&team
->barrier
))
2459 if (thr
->task
->taskgroup
)
2461 if (thr
->task
->taskgroup
->cancelled
)
2463 if (thr
->task
->taskgroup
->workshare
2464 && thr
->task
->taskgroup
->prev
2465 && thr
->task
->taskgroup
->prev
->cancelled
)
2470 /* The variables are mapped separately such that they can be released
2473 if ((flags
& GOMP_TARGET_FLAG_EXIT_DATA
) == 0)
2474 for (i
= 0; i
< mapnum
; i
++)
2475 if ((kinds
[i
] & 0xff) == GOMP_MAP_STRUCT
)
2477 gomp_map_vars (devicep
, sizes
[i
] + 1, &hostaddrs
[i
], NULL
, &sizes
[i
],
2478 &kinds
[i
], true, GOMP_MAP_VARS_ENTER_DATA
);
2481 else if ((kinds
[i
] & 0xff) == GOMP_MAP_TO_PSET
)
2483 for (j
= i
+ 1; j
< mapnum
; j
++)
2484 if (!GOMP_MAP_POINTER_P (get_kind (true, kinds
, j
) & 0xff))
2486 gomp_map_vars (devicep
, j
-i
, &hostaddrs
[i
], NULL
, &sizes
[i
],
2487 &kinds
[i
], true, GOMP_MAP_VARS_ENTER_DATA
);
2491 gomp_map_vars (devicep
, 1, &hostaddrs
[i
], NULL
, &sizes
[i
], &kinds
[i
],
2492 true, GOMP_MAP_VARS_ENTER_DATA
);
2494 gomp_exit_data (devicep
, mapnum
, hostaddrs
, sizes
, kinds
);
2498 gomp_target_task_fn (void *data
)
2500 struct gomp_target_task
*ttask
= (struct gomp_target_task
*) data
;
2501 struct gomp_device_descr
*devicep
= ttask
->devicep
;
2503 if (ttask
->fn
!= NULL
)
2507 || !(devicep
->capabilities
& GOMP_OFFLOAD_CAP_OPENMP_400
)
2508 || !(fn_addr
= gomp_get_target_fn_addr (devicep
, ttask
->fn
))
2509 || (devicep
->can_run_func
&& !devicep
->can_run_func (fn_addr
)))
2511 ttask
->state
= GOMP_TARGET_TASK_FALLBACK
;
2512 gomp_target_fallback (ttask
->fn
, ttask
->hostaddrs
);
2516 if (ttask
->state
== GOMP_TARGET_TASK_FINISHED
)
2519 gomp_unmap_vars (ttask
->tgt
, true);
2523 void *actual_arguments
;
2524 if (devicep
->capabilities
& GOMP_OFFLOAD_CAP_SHARED_MEM
)
2527 actual_arguments
= ttask
->hostaddrs
;
2531 ttask
->tgt
= gomp_map_vars (devicep
, ttask
->mapnum
, ttask
->hostaddrs
,
2532 NULL
, ttask
->sizes
, ttask
->kinds
, true,
2533 GOMP_MAP_VARS_TARGET
);
2534 actual_arguments
= (void *) ttask
->tgt
->tgt_start
;
2536 ttask
->state
= GOMP_TARGET_TASK_READY_TO_RUN
;
2538 assert (devicep
->async_run_func
);
2539 devicep
->async_run_func (devicep
->target_id
, fn_addr
, actual_arguments
,
2540 ttask
->args
, (void *) ttask
);
2543 else if (devicep
== NULL
2544 || !(devicep
->capabilities
& GOMP_OFFLOAD_CAP_OPENMP_400
)
2545 || devicep
->capabilities
& GOMP_OFFLOAD_CAP_SHARED_MEM
)
2549 if (ttask
->flags
& GOMP_TARGET_FLAG_UPDATE
)
2550 gomp_update (devicep
, ttask
->mapnum
, ttask
->hostaddrs
, ttask
->sizes
,
2551 ttask
->kinds
, true);
2552 else if ((ttask
->flags
& GOMP_TARGET_FLAG_EXIT_DATA
) == 0)
2553 for (i
= 0; i
< ttask
->mapnum
; i
++)
2554 if ((ttask
->kinds
[i
] & 0xff) == GOMP_MAP_STRUCT
)
2556 gomp_map_vars (devicep
, ttask
->sizes
[i
] + 1, &ttask
->hostaddrs
[i
],
2557 NULL
, &ttask
->sizes
[i
], &ttask
->kinds
[i
], true,
2558 GOMP_MAP_VARS_ENTER_DATA
);
2559 i
+= ttask
->sizes
[i
];
2562 gomp_map_vars (devicep
, 1, &ttask
->hostaddrs
[i
], NULL
, &ttask
->sizes
[i
],
2563 &ttask
->kinds
[i
], true, GOMP_MAP_VARS_ENTER_DATA
);
2565 gomp_exit_data (devicep
, ttask
->mapnum
, ttask
->hostaddrs
, ttask
->sizes
,
2571 GOMP_teams (unsigned int num_teams
, unsigned int thread_limit
)
2575 struct gomp_task_icv
*icv
= gomp_icv (true);
2576 icv
->thread_limit_var
2577 = thread_limit
> INT_MAX
? UINT_MAX
: thread_limit
;
2583 omp_target_alloc (size_t size
, int device_num
)
2585 if (device_num
== GOMP_DEVICE_HOST_FALLBACK
)
2586 return malloc (size
);
2591 struct gomp_device_descr
*devicep
= resolve_device (device_num
);
2592 if (devicep
== NULL
)
2595 if (!(devicep
->capabilities
& GOMP_OFFLOAD_CAP_OPENMP_400
)
2596 || devicep
->capabilities
& GOMP_OFFLOAD_CAP_SHARED_MEM
)
2597 return malloc (size
);
2599 gomp_mutex_lock (&devicep
->lock
);
2600 void *ret
= devicep
->alloc_func (devicep
->target_id
, size
);
2601 gomp_mutex_unlock (&devicep
->lock
);
2606 omp_target_free (void *device_ptr
, int device_num
)
2608 if (device_ptr
== NULL
)
2611 if (device_num
== GOMP_DEVICE_HOST_FALLBACK
)
2620 struct gomp_device_descr
*devicep
= resolve_device (device_num
);
2621 if (devicep
== NULL
)
2624 if (!(devicep
->capabilities
& GOMP_OFFLOAD_CAP_OPENMP_400
)
2625 || devicep
->capabilities
& GOMP_OFFLOAD_CAP_SHARED_MEM
)
2631 gomp_mutex_lock (&devicep
->lock
);
2632 gomp_free_device_memory (devicep
, device_ptr
);
2633 gomp_mutex_unlock (&devicep
->lock
);
2637 omp_target_is_present (const void *ptr
, int device_num
)
2642 if (device_num
== GOMP_DEVICE_HOST_FALLBACK
)
2648 struct gomp_device_descr
*devicep
= resolve_device (device_num
);
2649 if (devicep
== NULL
)
2652 if (!(devicep
->capabilities
& GOMP_OFFLOAD_CAP_OPENMP_400
)
2653 || devicep
->capabilities
& GOMP_OFFLOAD_CAP_SHARED_MEM
)
2656 gomp_mutex_lock (&devicep
->lock
);
2657 struct splay_tree_s
*mem_map
= &devicep
->mem_map
;
2658 struct splay_tree_key_s cur_node
;
2660 cur_node
.host_start
= (uintptr_t) ptr
;
2661 cur_node
.host_end
= cur_node
.host_start
;
2662 splay_tree_key n
= gomp_map_0len_lookup (mem_map
, &cur_node
);
2663 int ret
= n
!= NULL
;
2664 gomp_mutex_unlock (&devicep
->lock
);
2669 omp_target_memcpy (void *dst
, const void *src
, size_t length
,
2670 size_t dst_offset
, size_t src_offset
, int dst_device_num
,
2673 struct gomp_device_descr
*dst_devicep
= NULL
, *src_devicep
= NULL
;
2676 if (dst_device_num
!= GOMP_DEVICE_HOST_FALLBACK
)
2678 if (dst_device_num
< 0)
2681 dst_devicep
= resolve_device (dst_device_num
);
2682 if (dst_devicep
== NULL
)
2685 if (!(dst_devicep
->capabilities
& GOMP_OFFLOAD_CAP_OPENMP_400
)
2686 || dst_devicep
->capabilities
& GOMP_OFFLOAD_CAP_SHARED_MEM
)
2689 if (src_device_num
!= GOMP_DEVICE_HOST_FALLBACK
)
2691 if (src_device_num
< 0)
2694 src_devicep
= resolve_device (src_device_num
);
2695 if (src_devicep
== NULL
)
2698 if (!(src_devicep
->capabilities
& GOMP_OFFLOAD_CAP_OPENMP_400
)
2699 || src_devicep
->capabilities
& GOMP_OFFLOAD_CAP_SHARED_MEM
)
2702 if (src_devicep
== NULL
&& dst_devicep
== NULL
)
2704 memcpy ((char *) dst
+ dst_offset
, (char *) src
+ src_offset
, length
);
2707 if (src_devicep
== NULL
)
2709 gomp_mutex_lock (&dst_devicep
->lock
);
2710 ret
= dst_devicep
->host2dev_func (dst_devicep
->target_id
,
2711 (char *) dst
+ dst_offset
,
2712 (char *) src
+ src_offset
, length
);
2713 gomp_mutex_unlock (&dst_devicep
->lock
);
2714 return (ret
? 0 : EINVAL
);
2716 if (dst_devicep
== NULL
)
2718 gomp_mutex_lock (&src_devicep
->lock
);
2719 ret
= src_devicep
->dev2host_func (src_devicep
->target_id
,
2720 (char *) dst
+ dst_offset
,
2721 (char *) src
+ src_offset
, length
);
2722 gomp_mutex_unlock (&src_devicep
->lock
);
2723 return (ret
? 0 : EINVAL
);
2725 if (src_devicep
== dst_devicep
)
2727 gomp_mutex_lock (&src_devicep
->lock
);
2728 ret
= src_devicep
->dev2dev_func (src_devicep
->target_id
,
2729 (char *) dst
+ dst_offset
,
2730 (char *) src
+ src_offset
, length
);
2731 gomp_mutex_unlock (&src_devicep
->lock
);
2732 return (ret
? 0 : EINVAL
);
2738 omp_target_memcpy_rect_worker (void *dst
, const void *src
, size_t element_size
,
2739 int num_dims
, const size_t *volume
,
2740 const size_t *dst_offsets
,
2741 const size_t *src_offsets
,
2742 const size_t *dst_dimensions
,
2743 const size_t *src_dimensions
,
2744 struct gomp_device_descr
*dst_devicep
,
2745 struct gomp_device_descr
*src_devicep
)
2747 size_t dst_slice
= element_size
;
2748 size_t src_slice
= element_size
;
2749 size_t j
, dst_off
, src_off
, length
;
2754 if (__builtin_mul_overflow (element_size
, volume
[0], &length
)
2755 || __builtin_mul_overflow (element_size
, dst_offsets
[0], &dst_off
)
2756 || __builtin_mul_overflow (element_size
, src_offsets
[0], &src_off
))
2758 if (dst_devicep
== NULL
&& src_devicep
== NULL
)
2760 memcpy ((char *) dst
+ dst_off
, (const char *) src
+ src_off
,
2764 else if (src_devicep
== NULL
)
2765 ret
= dst_devicep
->host2dev_func (dst_devicep
->target_id
,
2766 (char *) dst
+ dst_off
,
2767 (const char *) src
+ src_off
,
2769 else if (dst_devicep
== NULL
)
2770 ret
= src_devicep
->dev2host_func (src_devicep
->target_id
,
2771 (char *) dst
+ dst_off
,
2772 (const char *) src
+ src_off
,
2774 else if (src_devicep
== dst_devicep
)
2775 ret
= src_devicep
->dev2dev_func (src_devicep
->target_id
,
2776 (char *) dst
+ dst_off
,
2777 (const char *) src
+ src_off
,
2781 return ret
? 0 : EINVAL
;
2784 /* FIXME: it would be nice to have some plugin function to handle
2785 num_dims == 2 and num_dims == 3 more efficiently. Larger ones can
2786 be handled in the generic recursion below, and for host-host it
2787 should be used even for any num_dims >= 2. */
2789 for (i
= 1; i
< num_dims
; i
++)
2790 if (__builtin_mul_overflow (dst_slice
, dst_dimensions
[i
], &dst_slice
)
2791 || __builtin_mul_overflow (src_slice
, src_dimensions
[i
], &src_slice
))
2793 if (__builtin_mul_overflow (dst_slice
, dst_offsets
[0], &dst_off
)
2794 || __builtin_mul_overflow (src_slice
, src_offsets
[0], &src_off
))
2796 for (j
= 0; j
< volume
[0]; j
++)
2798 ret
= omp_target_memcpy_rect_worker ((char *) dst
+ dst_off
,
2799 (const char *) src
+ src_off
,
2800 element_size
, num_dims
- 1,
2801 volume
+ 1, dst_offsets
+ 1,
2802 src_offsets
+ 1, dst_dimensions
+ 1,
2803 src_dimensions
+ 1, dst_devicep
,
2807 dst_off
+= dst_slice
;
2808 src_off
+= src_slice
;
2814 omp_target_memcpy_rect (void *dst
, const void *src
, size_t element_size
,
2815 int num_dims
, const size_t *volume
,
2816 const size_t *dst_offsets
,
2817 const size_t *src_offsets
,
2818 const size_t *dst_dimensions
,
2819 const size_t *src_dimensions
,
2820 int dst_device_num
, int src_device_num
)
2822 struct gomp_device_descr
*dst_devicep
= NULL
, *src_devicep
= NULL
;
2827 if (dst_device_num
!= GOMP_DEVICE_HOST_FALLBACK
)
2829 if (dst_device_num
< 0)
2832 dst_devicep
= resolve_device (dst_device_num
);
2833 if (dst_devicep
== NULL
)
2836 if (!(dst_devicep
->capabilities
& GOMP_OFFLOAD_CAP_OPENMP_400
)
2837 || dst_devicep
->capabilities
& GOMP_OFFLOAD_CAP_SHARED_MEM
)
2840 if (src_device_num
!= GOMP_DEVICE_HOST_FALLBACK
)
2842 if (src_device_num
< 0)
2845 src_devicep
= resolve_device (src_device_num
);
2846 if (src_devicep
== NULL
)
2849 if (!(src_devicep
->capabilities
& GOMP_OFFLOAD_CAP_OPENMP_400
)
2850 || src_devicep
->capabilities
& GOMP_OFFLOAD_CAP_SHARED_MEM
)
2854 if (src_devicep
!= NULL
&& dst_devicep
!= NULL
&& src_devicep
!= dst_devicep
)
2858 gomp_mutex_lock (&src_devicep
->lock
);
2859 else if (dst_devicep
)
2860 gomp_mutex_lock (&dst_devicep
->lock
);
2861 int ret
= omp_target_memcpy_rect_worker (dst
, src
, element_size
, num_dims
,
2862 volume
, dst_offsets
, src_offsets
,
2863 dst_dimensions
, src_dimensions
,
2864 dst_devicep
, src_devicep
);
2866 gomp_mutex_unlock (&src_devicep
->lock
);
2867 else if (dst_devicep
)
2868 gomp_mutex_unlock (&dst_devicep
->lock
);
2873 omp_target_associate_ptr (const void *host_ptr
, const void *device_ptr
,
2874 size_t size
, size_t device_offset
, int device_num
)
2876 if (device_num
== GOMP_DEVICE_HOST_FALLBACK
)
2882 struct gomp_device_descr
*devicep
= resolve_device (device_num
);
2883 if (devicep
== NULL
)
2886 if (!(devicep
->capabilities
& GOMP_OFFLOAD_CAP_OPENMP_400
)
2887 || devicep
->capabilities
& GOMP_OFFLOAD_CAP_SHARED_MEM
)
2890 gomp_mutex_lock (&devicep
->lock
);
2892 struct splay_tree_s
*mem_map
= &devicep
->mem_map
;
2893 struct splay_tree_key_s cur_node
;
2896 cur_node
.host_start
= (uintptr_t) host_ptr
;
2897 cur_node
.host_end
= cur_node
.host_start
+ size
;
2898 splay_tree_key n
= gomp_map_lookup (mem_map
, &cur_node
);
2901 if (n
->tgt
->tgt_start
+ n
->tgt_offset
2902 == (uintptr_t) device_ptr
+ device_offset
2903 && n
->host_start
<= cur_node
.host_start
2904 && n
->host_end
>= cur_node
.host_end
)
2909 struct target_mem_desc
*tgt
= gomp_malloc (sizeof (*tgt
));
2910 tgt
->array
= gomp_malloc (sizeof (*tgt
->array
));
2914 tgt
->to_free
= NULL
;
2916 tgt
->list_count
= 0;
2917 tgt
->device_descr
= devicep
;
2918 splay_tree_node array
= tgt
->array
;
2919 splay_tree_key k
= &array
->key
;
2920 k
->host_start
= cur_node
.host_start
;
2921 k
->host_end
= cur_node
.host_end
;
2923 k
->tgt_offset
= (uintptr_t) device_ptr
+ device_offset
;
2924 k
->refcount
= REFCOUNT_INFINITY
;
2925 k
->dynamic_refcount
= 0;
2928 array
->right
= NULL
;
2929 splay_tree_insert (&devicep
->mem_map
, array
);
2932 gomp_mutex_unlock (&devicep
->lock
);
2937 omp_target_disassociate_ptr (const void *ptr
, int device_num
)
2939 if (device_num
== GOMP_DEVICE_HOST_FALLBACK
)
2945 struct gomp_device_descr
*devicep
= resolve_device (device_num
);
2946 if (devicep
== NULL
)
2949 if (!(devicep
->capabilities
& GOMP_OFFLOAD_CAP_OPENMP_400
))
2952 gomp_mutex_lock (&devicep
->lock
);
2954 struct splay_tree_s
*mem_map
= &devicep
->mem_map
;
2955 struct splay_tree_key_s cur_node
;
2958 cur_node
.host_start
= (uintptr_t) ptr
;
2959 cur_node
.host_end
= cur_node
.host_start
;
2960 splay_tree_key n
= gomp_map_lookup (mem_map
, &cur_node
);
2962 && n
->host_start
== cur_node
.host_start
2963 && n
->refcount
== REFCOUNT_INFINITY
2964 && n
->tgt
->tgt_start
== 0
2965 && n
->tgt
->to_free
== NULL
2966 && n
->tgt
->refcount
== 1
2967 && n
->tgt
->list_count
== 0)
2969 splay_tree_remove (&devicep
->mem_map
, n
);
2970 gomp_unmap_tgt (n
->tgt
);
2974 gomp_mutex_unlock (&devicep
->lock
);
2979 omp_pause_resource (omp_pause_resource_t kind
, int device_num
)
2982 if (device_num
== GOMP_DEVICE_HOST_FALLBACK
)
2983 return gomp_pause_host ();
2984 if (device_num
< 0 || device_num
>= gomp_get_num_devices ())
2986 /* Do nothing for target devices for now. */
2991 omp_pause_resource_all (omp_pause_resource_t kind
)
2994 if (gomp_pause_host ())
2996 /* Do nothing for target devices for now. */
3000 ialias (omp_pause_resource
)
3001 ialias (omp_pause_resource_all
)
3003 #ifdef PLUGIN_SUPPORT
3005 /* This function tries to load a plugin for DEVICE. Name of plugin is passed
3007 The handles of the found functions are stored in the corresponding fields
3008 of DEVICE. The function returns TRUE on success and FALSE otherwise. */
3011 gomp_load_plugin_for_device (struct gomp_device_descr
*device
,
3012 const char *plugin_name
)
3014 const char *err
= NULL
, *last_missing
= NULL
;
3016 void *plugin_handle
= dlopen (plugin_name
, RTLD_LAZY
);
3020 /* Check if all required functions are available in the plugin and store
3021 their handlers. None of the symbols can legitimately be NULL,
3022 so we don't need to check dlerror all the time. */
3024 if (!(device->f##_func = dlsym (plugin_handle, "GOMP_OFFLOAD_" #f))) \
3026 /* Similar, but missing functions are not an error. Return false if
3027 failed, true otherwise. */
3028 #define DLSYM_OPT(f, n) \
3029 ((device->f##_func = dlsym (plugin_handle, "GOMP_OFFLOAD_" #n)) \
3030 || (last_missing = #n, 0))
3033 if (device
->version_func () != GOMP_VERSION
)
3035 err
= "plugin version mismatch";
3042 DLSYM (get_num_devices
);
3043 DLSYM (init_device
);
3044 DLSYM (fini_device
);
3046 DLSYM (unload_image
);
3051 device
->capabilities
= device
->get_caps_func ();
3052 if (device
->capabilities
& GOMP_OFFLOAD_CAP_OPENMP_400
)
3055 DLSYM_OPT (async_run
, async_run
);
3056 DLSYM_OPT (can_run
, can_run
);
3059 if (device
->capabilities
& GOMP_OFFLOAD_CAP_OPENACC_200
)
3061 if (!DLSYM_OPT (openacc
.exec
, openacc_exec
)
3062 || !DLSYM_OPT (openacc
.create_thread_data
,
3063 openacc_create_thread_data
)
3064 || !DLSYM_OPT (openacc
.destroy_thread_data
,
3065 openacc_destroy_thread_data
)
3066 || !DLSYM_OPT (openacc
.async
.construct
, openacc_async_construct
)
3067 || !DLSYM_OPT (openacc
.async
.destruct
, openacc_async_destruct
)
3068 || !DLSYM_OPT (openacc
.async
.test
, openacc_async_test
)
3069 || !DLSYM_OPT (openacc
.async
.synchronize
, openacc_async_synchronize
)
3070 || !DLSYM_OPT (openacc
.async
.serialize
, openacc_async_serialize
)
3071 || !DLSYM_OPT (openacc
.async
.queue_callback
,
3072 openacc_async_queue_callback
)
3073 || !DLSYM_OPT (openacc
.async
.exec
, openacc_async_exec
)
3074 || !DLSYM_OPT (openacc
.async
.dev2host
, openacc_async_dev2host
)
3075 || !DLSYM_OPT (openacc
.async
.host2dev
, openacc_async_host2dev
)
3076 || !DLSYM_OPT (openacc
.get_property
, openacc_get_property
))
3078 /* Require all the OpenACC handlers if we have
3079 GOMP_OFFLOAD_CAP_OPENACC_200. */
3080 err
= "plugin missing OpenACC handler function";
3085 cuda
+= DLSYM_OPT (openacc
.cuda
.get_current_device
,
3086 openacc_cuda_get_current_device
);
3087 cuda
+= DLSYM_OPT (openacc
.cuda
.get_current_context
,
3088 openacc_cuda_get_current_context
);
3089 cuda
+= DLSYM_OPT (openacc
.cuda
.get_stream
, openacc_cuda_get_stream
);
3090 cuda
+= DLSYM_OPT (openacc
.cuda
.set_stream
, openacc_cuda_set_stream
);
3091 if (cuda
&& cuda
!= 4)
3093 /* Make sure all the CUDA functions are there if any of them are. */
3094 err
= "plugin missing OpenACC CUDA handler function";
3106 gomp_error ("while loading %s: %s", plugin_name
, err
);
3108 gomp_error ("missing function was %s", last_missing
);
3110 dlclose (plugin_handle
);
3115 /* This function finalizes all initialized devices. */
3118 gomp_target_fini (void)
3121 for (i
= 0; i
< num_devices
; i
++)
3124 struct gomp_device_descr
*devicep
= &devices
[i
];
3125 gomp_mutex_lock (&devicep
->lock
);
3126 if (devicep
->state
== GOMP_DEVICE_INITIALIZED
)
3127 ret
= gomp_fini_device (devicep
);
3128 gomp_mutex_unlock (&devicep
->lock
);
3130 gomp_fatal ("device finalization failed");
3134 /* This function initializes the runtime for offloading.
3135 It parses the list of offload plugins, and tries to load these.
3136 On return, the variables NUM_DEVICES and NUM_DEVICES_OPENMP
3137 will be set, and the array DEVICES initialized, containing descriptors for
3138 corresponding devices, first the GOMP_OFFLOAD_CAP_OPENMP_400 ones, follows
3142 gomp_target_init (void)
3144 const char *prefix
="libgomp-plugin-";
3145 const char *suffix
= SONAME_SUFFIX (1);
3146 const char *cur
, *next
;
3148 int i
, new_num_devices
;
3153 cur
= OFFLOAD_PLUGINS
;
3157 struct gomp_device_descr current_device
;
3158 size_t prefix_len
, suffix_len
, cur_len
;
3160 next
= strchr (cur
, ',');
3162 prefix_len
= strlen (prefix
);
3163 cur_len
= next
? next
- cur
: strlen (cur
);
3164 suffix_len
= strlen (suffix
);
3166 plugin_name
= (char *) malloc (prefix_len
+ cur_len
+ suffix_len
+ 1);
3173 memcpy (plugin_name
, prefix
, prefix_len
);
3174 memcpy (plugin_name
+ prefix_len
, cur
, cur_len
);
3175 memcpy (plugin_name
+ prefix_len
+ cur_len
, suffix
, suffix_len
+ 1);
3177 if (gomp_load_plugin_for_device (¤t_device
, plugin_name
))
3179 new_num_devices
= current_device
.get_num_devices_func ();
3180 if (new_num_devices
>= 1)
3182 /* Augment DEVICES and NUM_DEVICES. */
3184 devices
= realloc (devices
, (num_devices
+ new_num_devices
)
3185 * sizeof (struct gomp_device_descr
));
3193 current_device
.name
= current_device
.get_name_func ();
3194 /* current_device.capabilities has already been set. */
3195 current_device
.type
= current_device
.get_type_func ();
3196 current_device
.mem_map
.root
= NULL
;
3197 current_device
.state
= GOMP_DEVICE_UNINITIALIZED
;
3198 for (i
= 0; i
< new_num_devices
; i
++)
3200 current_device
.target_id
= i
;
3201 devices
[num_devices
] = current_device
;
3202 gomp_mutex_init (&devices
[num_devices
].lock
);
3213 /* In DEVICES, sort the GOMP_OFFLOAD_CAP_OPENMP_400 ones first, and set
3214 NUM_DEVICES_OPENMP. */
3215 struct gomp_device_descr
*devices_s
3216 = malloc (num_devices
* sizeof (struct gomp_device_descr
));
3223 num_devices_openmp
= 0;
3224 for (i
= 0; i
< num_devices
; i
++)
3225 if (devices
[i
].capabilities
& GOMP_OFFLOAD_CAP_OPENMP_400
)
3226 devices_s
[num_devices_openmp
++] = devices
[i
];
3227 int num_devices_after_openmp
= num_devices_openmp
;
3228 for (i
= 0; i
< num_devices
; i
++)
3229 if (!(devices
[i
].capabilities
& GOMP_OFFLOAD_CAP_OPENMP_400
))
3230 devices_s
[num_devices_after_openmp
++] = devices
[i
];
3232 devices
= devices_s
;
3234 for (i
= 0; i
< num_devices
; i
++)
3236 /* The 'devices' array can be moved (by the realloc call) until we have
3237 found all the plugins, so registering with the OpenACC runtime (which
3238 takes a copy of the pointer argument) must be delayed until now. */
3239 if (devices
[i
].capabilities
& GOMP_OFFLOAD_CAP_OPENACC_200
)
3240 goacc_register (&devices
[i
]);
3243 if (atexit (gomp_target_fini
) != 0)
3244 gomp_fatal ("atexit failed");
3247 #else /* PLUGIN_SUPPORT */
3248 /* If dlfcn.h is unavailable we always fallback to host execution.
3249 GOMP_target* routines are just stubs for this case. */
3251 gomp_target_init (void)
3254 #endif /* PLUGIN_SUPPORT */