|
111
|
1 /* Copyright (C) 2013-2017 Free Software Foundation, Inc.
|
|
|
2 Contributed by Jakub Jelinek <jakub@redhat.com>.
|
|
|
3
|
|
|
4 This file is part of the GNU Offloading and Multi Processing Library
|
|
|
5 (libgomp).
|
|
|
6
|
|
|
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)
|
|
|
10 any later version.
|
|
|
11
|
|
|
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
|
|
|
15 more details.
|
|
|
16
|
|
|
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.
|
|
|
20
|
|
|
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/>. */
|
|
|
25
|
|
|
26 /* This file contains the support of offloading. */
|
|
|
27
|
|
|
28 #include "config.h"
|
|
|
29 #include "libgomp.h"
|
|
|
30 #include "oacc-plugin.h"
|
|
|
31 #include "oacc-int.h"
|
|
|
32 #include "gomp-constants.h"
|
|
|
33 #include <limits.h>
|
|
|
34 #include <stdbool.h>
|
|
|
35 #include <stdlib.h>
|
|
|
36 #ifdef HAVE_INTTYPES_H
|
|
|
37 # include <inttypes.h> /* For PRIu64. */
|
|
|
38 #endif
|
|
|
39 #include <string.h>
|
|
|
40 #include <assert.h>
|
|
|
41 #include <errno.h>
|
|
|
42
|
|
|
43 #ifdef PLUGIN_SUPPORT
|
|
|
44 #include <dlfcn.h>
|
|
|
45 #include "plugin-suffix.h"
|
|
|
46 #endif
|
|
|
47
|
|
|
48 static void gomp_target_init (void);
|
|
|
49
|
|
|
50 /* The whole initialization code for offloading plugins is only run one. */
|
|
|
51 static pthread_once_t gomp_is_initialized = PTHREAD_ONCE_INIT;
|
|
|
52
|
|
|
53 /* Mutex for offload image registration. */
|
|
|
54 static gomp_mutex_t register_lock;
|
|
|
55
|
|
|
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 {
|
|
|
60 unsigned version;
|
|
|
61 enum offload_target_type type;
|
|
|
62 const void *host_table;
|
|
|
63 const void *target_data;
|
|
|
64 };
|
|
|
65
|
|
|
66 /* Array of descriptors of offload images. */
|
|
|
67 static struct offload_image_descr *offload_images;
|
|
|
68
|
|
|
69 /* Total number of offload images. */
|
|
|
70 static int num_offload_images;
|
|
|
71
|
|
|
72 /* Array of descriptors for all available devices. */
|
|
|
73 static struct gomp_device_descr *devices;
|
|
|
74
|
|
|
75 /* Total number of available devices. */
|
|
|
76 static int num_devices;
|
|
|
77
|
|
|
78 /* Number of GOMP_OFFLOAD_CAP_OPENMP_400 devices. */
|
|
|
79 static int num_devices_openmp;
|
|
|
80
|
|
|
81 /* Similar to gomp_realloc, but release register_lock before gomp_fatal. */
|
|
|
82
|
|
|
83 static void *
|
|
|
84 gomp_realloc_unlock (void *old, size_t size)
|
|
|
85 {
|
|
|
86 void *ret = realloc (old, size);
|
|
|
87 if (ret == NULL)
|
|
|
88 {
|
|
|
89 gomp_mutex_unlock (®ister_lock);
|
|
|
90 gomp_fatal ("Out of memory allocating %lu bytes", (unsigned long) size);
|
|
|
91 }
|
|
|
92 return ret;
|
|
|
93 }
|
|
|
94
|
|
|
95 attribute_hidden void
|
|
|
96 gomp_init_targets_once (void)
|
|
|
97 {
|
|
|
98 (void) pthread_once (&gomp_is_initialized, gomp_target_init);
|
|
|
99 }
|
|
|
100
|
|
|
101 attribute_hidden int
|
|
|
102 gomp_get_num_devices (void)
|
|
|
103 {
|
|
|
104 gomp_init_targets_once ();
|
|
|
105 return num_devices_openmp;
|
|
|
106 }
|
|
|
107
|
|
|
108 static struct gomp_device_descr *
|
|
|
109 resolve_device (int device_id)
|
|
|
110 {
|
|
|
111 if (device_id == GOMP_DEVICE_ICV)
|
|
|
112 {
|
|
|
113 struct gomp_task_icv *icv = gomp_icv (false);
|
|
|
114 device_id = icv->default_device_var;
|
|
|
115 }
|
|
|
116
|
|
|
117 if (device_id < 0 || device_id >= gomp_get_num_devices ())
|
|
|
118 return NULL;
|
|
|
119
|
|
|
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)
|
|
|
124 {
|
|
|
125 gomp_mutex_unlock (&devices[device_id].lock);
|
|
|
126 return NULL;
|
|
|
127 }
|
|
|
128 gomp_mutex_unlock (&devices[device_id].lock);
|
|
|
129
|
|
|
130 return &devices[device_id];
|
|
|
131 }
|
|
|
132
|
|
|
133
|
|
|
134 static inline splay_tree_key
|
|
|
135 gomp_map_lookup (splay_tree mem_map, splay_tree_key key)
|
|
|
136 {
|
|
|
137 if (key->host_start != key->host_end)
|
|
|
138 return splay_tree_lookup (mem_map, key);
|
|
|
139
|
|
|
140 key->host_end++;
|
|
|
141 splay_tree_key n = splay_tree_lookup (mem_map, key);
|
|
|
142 key->host_end--;
|
|
|
143 if (n)
|
|
|
144 return n;
|
|
|
145 key->host_start--;
|
|
|
146 n = splay_tree_lookup (mem_map, key);
|
|
|
147 key->host_start++;
|
|
|
148 if (n)
|
|
|
149 return n;
|
|
|
150 return splay_tree_lookup (mem_map, key);
|
|
|
151 }
|
|
|
152
|
|
|
153 static inline splay_tree_key
|
|
|
154 gomp_map_0len_lookup (splay_tree mem_map, splay_tree_key key)
|
|
|
155 {
|
|
|
156 if (key->host_start != key->host_end)
|
|
|
157 return splay_tree_lookup (mem_map, key);
|
|
|
158
|
|
|
159 key->host_end++;
|
|
|
160 splay_tree_key n = splay_tree_lookup (mem_map, key);
|
|
|
161 key->host_end--;
|
|
|
162 return n;
|
|
|
163 }
|
|
|
164
|
|
|
165 static inline void
|
|
|
166 gomp_device_copy (struct gomp_device_descr *devicep,
|
|
|
167 bool (*copy_func) (int, void *, const void *, size_t),
|
|
|
168 const char *dst, void *dstaddr,
|
|
|
169 const char *src, const void *srcaddr,
|
|
|
170 size_t size)
|
|
|
171 {
|
|
|
172 if (!copy_func (devicep->target_id, dstaddr, srcaddr, size))
|
|
|
173 {
|
|
|
174 gomp_mutex_unlock (&devicep->lock);
|
|
|
175 gomp_fatal ("Copying of %s object [%p..%p) to %s object [%p..%p) failed",
|
|
|
176 src, srcaddr, srcaddr + size, dst, dstaddr, dstaddr + size);
|
|
|
177 }
|
|
|
178 }
|
|
|
179
|
|
|
180 static void
|
|
|
181 gomp_copy_host2dev (struct gomp_device_descr *devicep,
|
|
|
182 void *d, const void *h, size_t sz)
|
|
|
183 {
|
|
|
184 gomp_device_copy (devicep, devicep->host2dev_func, "dev", d, "host", h, sz);
|
|
|
185 }
|
|
|
186
|
|
|
187 static void
|
|
|
188 gomp_copy_dev2host (struct gomp_device_descr *devicep,
|
|
|
189 void *h, const void *d, size_t sz)
|
|
|
190 {
|
|
|
191 gomp_device_copy (devicep, devicep->dev2host_func, "host", h, "dev", d, sz);
|
|
|
192 }
|
|
|
193
|
|
|
194 static void
|
|
|
195 gomp_free_device_memory (struct gomp_device_descr *devicep, void *devptr)
|
|
|
196 {
|
|
|
197 if (!devicep->free_func (devicep->target_id, devptr))
|
|
|
198 {
|
|
|
199 gomp_mutex_unlock (&devicep->lock);
|
|
|
200 gomp_fatal ("error in freeing device memory block at %p", devptr);
|
|
|
201 }
|
|
|
202 }
|
|
|
203
|
|
|
204 /* Handle the case where gomp_map_lookup, splay_tree_lookup or
|
|
|
205 gomp_map_0len_lookup found oldn for newn.
|
|
|
206 Helper function of gomp_map_vars. */
|
|
|
207
|
|
|
208 static inline void
|
|
|
209 gomp_map_vars_existing (struct gomp_device_descr *devicep, splay_tree_key oldn,
|
|
|
210 splay_tree_key newn, struct target_var_desc *tgt_var,
|
|
|
211 unsigned char kind)
|
|
|
212 {
|
|
|
213 tgt_var->key = oldn;
|
|
|
214 tgt_var->copy_from = GOMP_MAP_COPY_FROM_P (kind);
|
|
|
215 tgt_var->always_copy_from = GOMP_MAP_ALWAYS_FROM_P (kind);
|
|
|
216 tgt_var->offset = newn->host_start - oldn->host_start;
|
|
|
217 tgt_var->length = newn->host_end - newn->host_start;
|
|
|
218
|
|
|
219 if ((kind & GOMP_MAP_FLAG_FORCE)
|
|
|
220 || oldn->host_start > newn->host_start
|
|
|
221 || oldn->host_end < newn->host_end)
|
|
|
222 {
|
|
|
223 gomp_mutex_unlock (&devicep->lock);
|
|
|
224 gomp_fatal ("Trying to map into device [%p..%p) object when "
|
|
|
225 "[%p..%p) is already mapped",
|
|
|
226 (void *) newn->host_start, (void *) newn->host_end,
|
|
|
227 (void *) oldn->host_start, (void *) oldn->host_end);
|
|
|
228 }
|
|
|
229
|
|
|
230 if (GOMP_MAP_ALWAYS_TO_P (kind))
|
|
|
231 gomp_copy_host2dev (devicep,
|
|
|
232 (void *) (oldn->tgt->tgt_start + oldn->tgt_offset
|
|
|
233 + newn->host_start - oldn->host_start),
|
|
|
234 (void *) newn->host_start,
|
|
|
235 newn->host_end - newn->host_start);
|
|
|
236
|
|
|
237 if (oldn->refcount != REFCOUNT_INFINITY)
|
|
|
238 oldn->refcount++;
|
|
|
239 }
|
|
|
240
|
|
|
241 static int
|
|
|
242 get_kind (bool short_mapkind, void *kinds, int idx)
|
|
|
243 {
|
|
|
244 return short_mapkind ? ((unsigned short *) kinds)[idx]
|
|
|
245 : ((unsigned char *) kinds)[idx];
|
|
|
246 }
|
|
|
247
|
|
|
248 static void
|
|
|
249 gomp_map_pointer (struct target_mem_desc *tgt, uintptr_t host_ptr,
|
|
|
250 uintptr_t target_offset, uintptr_t bias)
|
|
|
251 {
|
|
|
252 struct gomp_device_descr *devicep = tgt->device_descr;
|
|
|
253 struct splay_tree_s *mem_map = &devicep->mem_map;
|
|
|
254 struct splay_tree_key_s cur_node;
|
|
|
255
|
|
|
256 cur_node.host_start = host_ptr;
|
|
|
257 if (cur_node.host_start == (uintptr_t) NULL)
|
|
|
258 {
|
|
|
259 cur_node.tgt_offset = (uintptr_t) NULL;
|
|
|
260 /* FIXME: see comment about coalescing host/dev transfers below. */
|
|
|
261 gomp_copy_host2dev (devicep,
|
|
|
262 (void *) (tgt->tgt_start + target_offset),
|
|
|
263 (void *) &cur_node.tgt_offset,
|
|
|
264 sizeof (void *));
|
|
|
265 return;
|
|
|
266 }
|
|
|
267 /* Add bias to the pointer value. */
|
|
|
268 cur_node.host_start += bias;
|
|
|
269 cur_node.host_end = cur_node.host_start;
|
|
|
270 splay_tree_key n = gomp_map_lookup (mem_map, &cur_node);
|
|
|
271 if (n == NULL)
|
|
|
272 {
|
|
|
273 gomp_mutex_unlock (&devicep->lock);
|
|
|
274 gomp_fatal ("Pointer target of array section wasn't mapped");
|
|
|
275 }
|
|
|
276 cur_node.host_start -= n->host_start;
|
|
|
277 cur_node.tgt_offset
|
|
|
278 = n->tgt->tgt_start + n->tgt_offset + cur_node.host_start;
|
|
|
279 /* At this point tgt_offset is target address of the
|
|
|
280 array section. Now subtract bias to get what we want
|
|
|
281 to initialize the pointer with. */
|
|
|
282 cur_node.tgt_offset -= bias;
|
|
|
283 /* FIXME: see comment about coalescing host/dev transfers below. */
|
|
|
284 gomp_copy_host2dev (devicep, (void *) (tgt->tgt_start + target_offset),
|
|
|
285 (void *) &cur_node.tgt_offset, sizeof (void *));
|
|
|
286 }
|
|
|
287
|
|
|
288 static void
|
|
|
289 gomp_map_fields_existing (struct target_mem_desc *tgt, splay_tree_key n,
|
|
|
290 size_t first, size_t i, void **hostaddrs,
|
|
|
291 size_t *sizes, void *kinds)
|
|
|
292 {
|
|
|
293 struct gomp_device_descr *devicep = tgt->device_descr;
|
|
|
294 struct splay_tree_s *mem_map = &devicep->mem_map;
|
|
|
295 struct splay_tree_key_s cur_node;
|
|
|
296 int kind;
|
|
|
297 const bool short_mapkind = true;
|
|
|
298 const int typemask = short_mapkind ? 0xff : 0x7;
|
|
|
299
|
|
|
300 cur_node.host_start = (uintptr_t) hostaddrs[i];
|
|
|
301 cur_node.host_end = cur_node.host_start + sizes[i];
|
|
|
302 splay_tree_key n2 = splay_tree_lookup (mem_map, &cur_node);
|
|
|
303 kind = get_kind (short_mapkind, kinds, i);
|
|
|
304 if (n2
|
|
|
305 && n2->tgt == n->tgt
|
|
|
306 && n2->host_start - n->host_start == n2->tgt_offset - n->tgt_offset)
|
|
|
307 {
|
|
|
308 gomp_map_vars_existing (devicep, n2, &cur_node,
|
|
|
309 &tgt->list[i], kind & typemask);
|
|
|
310 return;
|
|
|
311 }
|
|
|
312 if (sizes[i] == 0)
|
|
|
313 {
|
|
|
314 if (cur_node.host_start > (uintptr_t) hostaddrs[first - 1])
|
|
|
315 {
|
|
|
316 cur_node.host_start--;
|
|
|
317 n2 = splay_tree_lookup (mem_map, &cur_node);
|
|
|
318 cur_node.host_start++;
|
|
|
319 if (n2
|
|
|
320 && n2->tgt == n->tgt
|
|
|
321 && n2->host_start - n->host_start
|
|
|
322 == n2->tgt_offset - n->tgt_offset)
|
|
|
323 {
|
|
|
324 gomp_map_vars_existing (devicep, n2, &cur_node, &tgt->list[i],
|
|
|
325 kind & typemask);
|
|
|
326 return;
|
|
|
327 }
|
|
|
328 }
|
|
|
329 cur_node.host_end++;
|
|
|
330 n2 = splay_tree_lookup (mem_map, &cur_node);
|
|
|
331 cur_node.host_end--;
|
|
|
332 if (n2
|
|
|
333 && n2->tgt == n->tgt
|
|
|
334 && n2->host_start - n->host_start == n2->tgt_offset - n->tgt_offset)
|
|
|
335 {
|
|
|
336 gomp_map_vars_existing (devicep, n2, &cur_node, &tgt->list[i],
|
|
|
337 kind & typemask);
|
|
|
338 return;
|
|
|
339 }
|
|
|
340 }
|
|
|
341 gomp_mutex_unlock (&devicep->lock);
|
|
|
342 gomp_fatal ("Trying to map into device [%p..%p) structure element when "
|
|
|
343 "other mapped elements from the same structure weren't mapped "
|
|
|
344 "together with it", (void *) cur_node.host_start,
|
|
|
345 (void *) cur_node.host_end);
|
|
|
346 }
|
|
|
347
|
|
|
348 static inline uintptr_t
|
|
|
349 gomp_map_val (struct target_mem_desc *tgt, void **hostaddrs, size_t i)
|
|
|
350 {
|
|
|
351 if (tgt->list[i].key != NULL)
|
|
|
352 return tgt->list[i].key->tgt->tgt_start
|
|
|
353 + tgt->list[i].key->tgt_offset
|
|
|
354 + tgt->list[i].offset;
|
|
|
355 if (tgt->list[i].offset == ~(uintptr_t) 0)
|
|
|
356 return (uintptr_t) hostaddrs[i];
|
|
|
357 if (tgt->list[i].offset == ~(uintptr_t) 1)
|
|
|
358 return 0;
|
|
|
359 if (tgt->list[i].offset == ~(uintptr_t) 2)
|
|
|
360 return tgt->list[i + 1].key->tgt->tgt_start
|
|
|
361 + tgt->list[i + 1].key->tgt_offset
|
|
|
362 + tgt->list[i + 1].offset
|
|
|
363 + (uintptr_t) hostaddrs[i]
|
|
|
364 - (uintptr_t) hostaddrs[i + 1];
|
|
|
365 return tgt->tgt_start + tgt->list[i].offset;
|
|
|
366 }
|
|
|
367
|
|
|
368 attribute_hidden struct target_mem_desc *
|
|
|
369 gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
|
|
|
370 void **hostaddrs, void **devaddrs, size_t *sizes, void *kinds,
|
|
|
371 bool short_mapkind, enum gomp_map_vars_kind pragma_kind)
|
|
|
372 {
|
|
|
373 size_t i, tgt_align, tgt_size, not_found_cnt = 0;
|
|
|
374 bool has_firstprivate = false;
|
|
|
375 const int rshift = short_mapkind ? 8 : 3;
|
|
|
376 const int typemask = short_mapkind ? 0xff : 0x7;
|
|
|
377 struct splay_tree_s *mem_map = &devicep->mem_map;
|
|
|
378 struct splay_tree_key_s cur_node;
|
|
|
379 struct target_mem_desc *tgt
|
|
|
380 = gomp_malloc (sizeof (*tgt) + sizeof (tgt->list[0]) * mapnum);
|
|
|
381 tgt->list_count = mapnum;
|
|
|
382 tgt->refcount = pragma_kind == GOMP_MAP_VARS_ENTER_DATA ? 0 : 1;
|
|
|
383 tgt->device_descr = devicep;
|
|
|
384
|
|
|
385 if (mapnum == 0)
|
|
|
386 {
|
|
|
387 tgt->tgt_start = 0;
|
|
|
388 tgt->tgt_end = 0;
|
|
|
389 return tgt;
|
|
|
390 }
|
|
|
391
|
|
|
392 tgt_align = sizeof (void *);
|
|
|
393 tgt_size = 0;
|
|
|
394 if (pragma_kind == GOMP_MAP_VARS_TARGET)
|
|
|
395 {
|
|
|
396 size_t align = 4 * sizeof (void *);
|
|
|
397 tgt_align = align;
|
|
|
398 tgt_size = mapnum * sizeof (void *);
|
|
|
399 }
|
|
|
400
|
|
|
401 gomp_mutex_lock (&devicep->lock);
|
|
|
402 if (devicep->state == GOMP_DEVICE_FINALIZED)
|
|
|
403 {
|
|
|
404 gomp_mutex_unlock (&devicep->lock);
|
|
|
405 free (tgt);
|
|
|
406 return NULL;
|
|
|
407 }
|
|
|
408
|
|
|
409 for (i = 0; i < mapnum; i++)
|
|
|
410 {
|
|
|
411 int kind = get_kind (short_mapkind, kinds, i);
|
|
|
412 if (hostaddrs[i] == NULL
|
|
|
413 || (kind & typemask) == GOMP_MAP_FIRSTPRIVATE_INT)
|
|
|
414 {
|
|
|
415 tgt->list[i].key = NULL;
|
|
|
416 tgt->list[i].offset = ~(uintptr_t) 0;
|
|
|
417 continue;
|
|
|
418 }
|
|
|
419 else if ((kind & typemask) == GOMP_MAP_USE_DEVICE_PTR)
|
|
|
420 {
|
|
|
421 cur_node.host_start = (uintptr_t) hostaddrs[i];
|
|
|
422 cur_node.host_end = cur_node.host_start;
|
|
|
423 splay_tree_key n = gomp_map_lookup (mem_map, &cur_node);
|
|
|
424 if (n == NULL)
|
|
|
425 {
|
|
|
426 gomp_mutex_unlock (&devicep->lock);
|
|
|
427 gomp_fatal ("use_device_ptr pointer wasn't mapped");
|
|
|
428 }
|
|
|
429 cur_node.host_start -= n->host_start;
|
|
|
430 hostaddrs[i]
|
|
|
431 = (void *) (n->tgt->tgt_start + n->tgt_offset
|
|
|
432 + cur_node.host_start);
|
|
|
433 tgt->list[i].key = NULL;
|
|
|
434 tgt->list[i].offset = ~(uintptr_t) 0;
|
|
|
435 continue;
|
|
|
436 }
|
|
|
437 else if ((kind & typemask) == GOMP_MAP_STRUCT)
|
|
|
438 {
|
|
|
439 size_t first = i + 1;
|
|
|
440 size_t last = i + sizes[i];
|
|
|
441 cur_node.host_start = (uintptr_t) hostaddrs[i];
|
|
|
442 cur_node.host_end = (uintptr_t) hostaddrs[last]
|
|
|
443 + sizes[last];
|
|
|
444 tgt->list[i].key = NULL;
|
|
|
445 tgt->list[i].offset = ~(uintptr_t) 2;
|
|
|
446 splay_tree_key n = splay_tree_lookup (mem_map, &cur_node);
|
|
|
447 if (n == NULL)
|
|
|
448 {
|
|
|
449 size_t align = (size_t) 1 << (kind >> rshift);
|
|
|
450 if (tgt_align < align)
|
|
|
451 tgt_align = align;
|
|
|
452 tgt_size -= (uintptr_t) hostaddrs[first]
|
|
|
453 - (uintptr_t) hostaddrs[i];
|
|
|
454 tgt_size = (tgt_size + align - 1) & ~(align - 1);
|
|
|
455 tgt_size += cur_node.host_end - (uintptr_t) hostaddrs[i];
|
|
|
456 not_found_cnt += last - i;
|
|
|
457 for (i = first; i <= last; i++)
|
|
|
458 tgt->list[i].key = NULL;
|
|
|
459 i--;
|
|
|
460 continue;
|
|
|
461 }
|
|
|
462 for (i = first; i <= last; i++)
|
|
|
463 gomp_map_fields_existing (tgt, n, first, i, hostaddrs,
|
|
|
464 sizes, kinds);
|
|
|
465 i--;
|
|
|
466 continue;
|
|
|
467 }
|
|
|
468 else if ((kind & typemask) == GOMP_MAP_ALWAYS_POINTER)
|
|
|
469 {
|
|
|
470 tgt->list[i].key = NULL;
|
|
|
471 tgt->list[i].offset = ~(uintptr_t) 1;
|
|
|
472 has_firstprivate = true;
|
|
|
473 continue;
|
|
|
474 }
|
|
|
475 cur_node.host_start = (uintptr_t) hostaddrs[i];
|
|
|
476 if (!GOMP_MAP_POINTER_P (kind & typemask))
|
|
|
477 cur_node.host_end = cur_node.host_start + sizes[i];
|
|
|
478 else
|
|
|
479 cur_node.host_end = cur_node.host_start + sizeof (void *);
|
|
|
480 if ((kind & typemask) == GOMP_MAP_FIRSTPRIVATE)
|
|
|
481 {
|
|
|
482 tgt->list[i].key = NULL;
|
|
|
483
|
|
|
484 size_t align = (size_t) 1 << (kind >> rshift);
|
|
|
485 if (tgt_align < align)
|
|
|
486 tgt_align = align;
|
|
|
487 tgt_size = (tgt_size + align - 1) & ~(align - 1);
|
|
|
488 tgt_size += cur_node.host_end - cur_node.host_start;
|
|
|
489 has_firstprivate = true;
|
|
|
490 continue;
|
|
|
491 }
|
|
|
492 splay_tree_key n;
|
|
|
493 if ((kind & typemask) == GOMP_MAP_ZERO_LEN_ARRAY_SECTION)
|
|
|
494 {
|
|
|
495 n = gomp_map_0len_lookup (mem_map, &cur_node);
|
|
|
496 if (!n)
|
|
|
497 {
|
|
|
498 tgt->list[i].key = NULL;
|
|
|
499 tgt->list[i].offset = ~(uintptr_t) 1;
|
|
|
500 continue;
|
|
|
501 }
|
|
|
502 }
|
|
|
503 else
|
|
|
504 n = splay_tree_lookup (mem_map, &cur_node);
|
|
|
505 if (n && n->refcount != REFCOUNT_LINK)
|
|
|
506 gomp_map_vars_existing (devicep, n, &cur_node, &tgt->list[i],
|
|
|
507 kind & typemask);
|
|
|
508 else
|
|
|
509 {
|
|
|
510 tgt->list[i].key = NULL;
|
|
|
511
|
|
|
512 size_t align = (size_t) 1 << (kind >> rshift);
|
|
|
513 not_found_cnt++;
|
|
|
514 if (tgt_align < align)
|
|
|
515 tgt_align = align;
|
|
|
516 tgt_size = (tgt_size + align - 1) & ~(align - 1);
|
|
|
517 tgt_size += cur_node.host_end - cur_node.host_start;
|
|
|
518 if ((kind & typemask) == GOMP_MAP_TO_PSET)
|
|
|
519 {
|
|
|
520 size_t j;
|
|
|
521 for (j = i + 1; j < mapnum; j++)
|
|
|
522 if (!GOMP_MAP_POINTER_P (get_kind (short_mapkind, kinds, j)
|
|
|
523 & typemask))
|
|
|
524 break;
|
|
|
525 else if ((uintptr_t) hostaddrs[j] < cur_node.host_start
|
|
|
526 || ((uintptr_t) hostaddrs[j] + sizeof (void *)
|
|
|
527 > cur_node.host_end))
|
|
|
528 break;
|
|
|
529 else
|
|
|
530 {
|
|
|
531 tgt->list[j].key = NULL;
|
|
|
532 i++;
|
|
|
533 }
|
|
|
534 }
|
|
|
535 }
|
|
|
536 }
|
|
|
537
|
|
|
538 if (devaddrs)
|
|
|
539 {
|
|
|
540 if (mapnum != 1)
|
|
|
541 {
|
|
|
542 gomp_mutex_unlock (&devicep->lock);
|
|
|
543 gomp_fatal ("unexpected aggregation");
|
|
|
544 }
|
|
|
545 tgt->to_free = devaddrs[0];
|
|
|
546 tgt->tgt_start = (uintptr_t) tgt->to_free;
|
|
|
547 tgt->tgt_end = tgt->tgt_start + sizes[0];
|
|
|
548 }
|
|
|
549 else if (not_found_cnt || pragma_kind == GOMP_MAP_VARS_TARGET)
|
|
|
550 {
|
|
|
551 /* Allocate tgt_align aligned tgt_size block of memory. */
|
|
|
552 /* FIXME: Perhaps change interface to allocate properly aligned
|
|
|
553 memory. */
|
|
|
554 tgt->to_free = devicep->alloc_func (devicep->target_id,
|
|
|
555 tgt_size + tgt_align - 1);
|
|
|
556 if (!tgt->to_free)
|
|
|
557 {
|
|
|
558 gomp_mutex_unlock (&devicep->lock);
|
|
|
559 gomp_fatal ("device memory allocation fail");
|
|
|
560 }
|
|
|
561
|
|
|
562 tgt->tgt_start = (uintptr_t) tgt->to_free;
|
|
|
563 tgt->tgt_start = (tgt->tgt_start + tgt_align - 1) & ~(tgt_align - 1);
|
|
|
564 tgt->tgt_end = tgt->tgt_start + tgt_size;
|
|
|
565 }
|
|
|
566 else
|
|
|
567 {
|
|
|
568 tgt->to_free = NULL;
|
|
|
569 tgt->tgt_start = 0;
|
|
|
570 tgt->tgt_end = 0;
|
|
|
571 }
|
|
|
572
|
|
|
573 tgt_size = 0;
|
|
|
574 if (pragma_kind == GOMP_MAP_VARS_TARGET)
|
|
|
575 tgt_size = mapnum * sizeof (void *);
|
|
|
576
|
|
|
577 tgt->array = NULL;
|
|
|
578 if (not_found_cnt || has_firstprivate)
|
|
|
579 {
|
|
|
580 if (not_found_cnt)
|
|
|
581 tgt->array = gomp_malloc (not_found_cnt * sizeof (*tgt->array));
|
|
|
582 splay_tree_node array = tgt->array;
|
|
|
583 size_t j, field_tgt_offset = 0, field_tgt_clear = ~(size_t) 0;
|
|
|
584 uintptr_t field_tgt_base = 0;
|
|
|
585
|
|
|
586 for (i = 0; i < mapnum; i++)
|
|
|
587 if (tgt->list[i].key == NULL)
|
|
|
588 {
|
|
|
589 int kind = get_kind (short_mapkind, kinds, i);
|
|
|
590 if (hostaddrs[i] == NULL)
|
|
|
591 continue;
|
|
|
592 switch (kind & typemask)
|
|
|
593 {
|
|
|
594 size_t align, len, first, last;
|
|
|
595 splay_tree_key n;
|
|
|
596 case GOMP_MAP_FIRSTPRIVATE:
|
|
|
597 align = (size_t) 1 << (kind >> rshift);
|
|
|
598 tgt_size = (tgt_size + align - 1) & ~(align - 1);
|
|
|
599 tgt->list[i].offset = tgt_size;
|
|
|
600 len = sizes[i];
|
|
|
601 gomp_copy_host2dev (devicep,
|
|
|
602 (void *) (tgt->tgt_start + tgt_size),
|
|
|
603 (void *) hostaddrs[i], len);
|
|
|
604 tgt_size += len;
|
|
|
605 continue;
|
|
|
606 case GOMP_MAP_FIRSTPRIVATE_INT:
|
|
|
607 case GOMP_MAP_USE_DEVICE_PTR:
|
|
|
608 case GOMP_MAP_ZERO_LEN_ARRAY_SECTION:
|
|
|
609 continue;
|
|
|
610 case GOMP_MAP_STRUCT:
|
|
|
611 first = i + 1;
|
|
|
612 last = i + sizes[i];
|
|
|
613 cur_node.host_start = (uintptr_t) hostaddrs[i];
|
|
|
614 cur_node.host_end = (uintptr_t) hostaddrs[last]
|
|
|
615 + sizes[last];
|
|
|
616 if (tgt->list[first].key != NULL)
|
|
|
617 continue;
|
|
|
618 n = splay_tree_lookup (mem_map, &cur_node);
|
|
|
619 if (n == NULL)
|
|
|
620 {
|
|
|
621 size_t align = (size_t) 1 << (kind >> rshift);
|
|
|
622 tgt_size -= (uintptr_t) hostaddrs[first]
|
|
|
623 - (uintptr_t) hostaddrs[i];
|
|
|
624 tgt_size = (tgt_size + align - 1) & ~(align - 1);
|
|
|
625 tgt_size += (uintptr_t) hostaddrs[first]
|
|
|
626 - (uintptr_t) hostaddrs[i];
|
|
|
627 field_tgt_base = (uintptr_t) hostaddrs[first];
|
|
|
628 field_tgt_offset = tgt_size;
|
|
|
629 field_tgt_clear = last;
|
|
|
630 tgt_size += cur_node.host_end
|
|
|
631 - (uintptr_t) hostaddrs[first];
|
|
|
632 continue;
|
|
|
633 }
|
|
|
634 for (i = first; i <= last; i++)
|
|
|
635 gomp_map_fields_existing (tgt, n, first, i, hostaddrs,
|
|
|
636 sizes, kinds);
|
|
|
637 i--;
|
|
|
638 continue;
|
|
|
639 case GOMP_MAP_ALWAYS_POINTER:
|
|
|
640 cur_node.host_start = (uintptr_t) hostaddrs[i];
|
|
|
641 cur_node.host_end = cur_node.host_start + sizeof (void *);
|
|
|
642 n = splay_tree_lookup (mem_map, &cur_node);
|
|
|
643 if (n == NULL
|
|
|
644 || n->host_start > cur_node.host_start
|
|
|
645 || n->host_end < cur_node.host_end)
|
|
|
646 {
|
|
|
647 gomp_mutex_unlock (&devicep->lock);
|
|
|
648 gomp_fatal ("always pointer not mapped");
|
|
|
649 }
|
|
|
650 if ((get_kind (short_mapkind, kinds, i - 1) & typemask)
|
|
|
651 != GOMP_MAP_ALWAYS_POINTER)
|
|
|
652 cur_node.tgt_offset = gomp_map_val (tgt, hostaddrs, i - 1);
|
|
|
653 if (cur_node.tgt_offset)
|
|
|
654 cur_node.tgt_offset -= sizes[i];
|
|
|
655 gomp_copy_host2dev (devicep,
|
|
|
656 (void *) (n->tgt->tgt_start
|
|
|
657 + n->tgt_offset
|
|
|
658 + cur_node.host_start
|
|
|
659 - n->host_start),
|
|
|
660 (void *) &cur_node.tgt_offset,
|
|
|
661 sizeof (void *));
|
|
|
662 cur_node.tgt_offset = n->tgt->tgt_start + n->tgt_offset
|
|
|
663 + cur_node.host_start - n->host_start;
|
|
|
664 continue;
|
|
|
665 default:
|
|
|
666 break;
|
|
|
667 }
|
|
|
668 splay_tree_key k = &array->key;
|
|
|
669 k->host_start = (uintptr_t) hostaddrs[i];
|
|
|
670 if (!GOMP_MAP_POINTER_P (kind & typemask))
|
|
|
671 k->host_end = k->host_start + sizes[i];
|
|
|
672 else
|
|
|
673 k->host_end = k->host_start + sizeof (void *);
|
|
|
674 splay_tree_key n = splay_tree_lookup (mem_map, k);
|
|
|
675 if (n && n->refcount != REFCOUNT_LINK)
|
|
|
676 gomp_map_vars_existing (devicep, n, k, &tgt->list[i],
|
|
|
677 kind & typemask);
|
|
|
678 else
|
|
|
679 {
|
|
|
680 k->link_key = NULL;
|
|
|
681 if (n && n->refcount == REFCOUNT_LINK)
|
|
|
682 {
|
|
|
683 /* Replace target address of the pointer with target address
|
|
|
684 of mapped object in the splay tree. */
|
|
|
685 splay_tree_remove (mem_map, n);
|
|
|
686 k->link_key = n;
|
|
|
687 }
|
|
|
688 size_t align = (size_t) 1 << (kind >> rshift);
|
|
|
689 tgt->list[i].key = k;
|
|
|
690 k->tgt = tgt;
|
|
|
691 if (field_tgt_clear != ~(size_t) 0)
|
|
|
692 {
|
|
|
693 k->tgt_offset = k->host_start - field_tgt_base
|
|
|
694 + field_tgt_offset;
|
|
|
695 if (i == field_tgt_clear)
|
|
|
696 field_tgt_clear = ~(size_t) 0;
|
|
|
697 }
|
|
|
698 else
|
|
|
699 {
|
|
|
700 tgt_size = (tgt_size + align - 1) & ~(align - 1);
|
|
|
701 k->tgt_offset = tgt_size;
|
|
|
702 tgt_size += k->host_end - k->host_start;
|
|
|
703 }
|
|
|
704 tgt->list[i].copy_from = GOMP_MAP_COPY_FROM_P (kind & typemask);
|
|
|
705 tgt->list[i].always_copy_from
|
|
|
706 = GOMP_MAP_ALWAYS_FROM_P (kind & typemask);
|
|
|
707 tgt->list[i].offset = 0;
|
|
|
708 tgt->list[i].length = k->host_end - k->host_start;
|
|
|
709 k->refcount = 1;
|
|
|
710 tgt->refcount++;
|
|
|
711 array->left = NULL;
|
|
|
712 array->right = NULL;
|
|
|
713 splay_tree_insert (mem_map, array);
|
|
|
714 switch (kind & typemask)
|
|
|
715 {
|
|
|
716 case GOMP_MAP_ALLOC:
|
|
|
717 case GOMP_MAP_FROM:
|
|
|
718 case GOMP_MAP_FORCE_ALLOC:
|
|
|
719 case GOMP_MAP_FORCE_FROM:
|
|
|
720 case GOMP_MAP_ALWAYS_FROM:
|
|
|
721 break;
|
|
|
722 case GOMP_MAP_TO:
|
|
|
723 case GOMP_MAP_TOFROM:
|
|
|
724 case GOMP_MAP_FORCE_TO:
|
|
|
725 case GOMP_MAP_FORCE_TOFROM:
|
|
|
726 case GOMP_MAP_ALWAYS_TO:
|
|
|
727 case GOMP_MAP_ALWAYS_TOFROM:
|
|
|
728 /* FIXME: Perhaps add some smarts, like if copying
|
|
|
729 several adjacent fields from host to target, use some
|
|
|
730 host buffer to avoid sending each var individually. */
|
|
|
731 gomp_copy_host2dev (devicep,
|
|
|
732 (void *) (tgt->tgt_start
|
|
|
733 + k->tgt_offset),
|
|
|
734 (void *) k->host_start,
|
|
|
735 k->host_end - k->host_start);
|
|
|
736 break;
|
|
|
737 case GOMP_MAP_POINTER:
|
|
|
738 gomp_map_pointer (tgt, (uintptr_t) *(void **) k->host_start,
|
|
|
739 k->tgt_offset, sizes[i]);
|
|
|
740 break;
|
|
|
741 case GOMP_MAP_TO_PSET:
|
|
|
742 /* FIXME: see above FIXME comment. */
|
|
|
743 gomp_copy_host2dev (devicep,
|
|
|
744 (void *) (tgt->tgt_start
|
|
|
745 + k->tgt_offset),
|
|
|
746 (void *) k->host_start,
|
|
|
747 k->host_end - k->host_start);
|
|
|
748
|
|
|
749 for (j = i + 1; j < mapnum; j++)
|
|
|
750 if (!GOMP_MAP_POINTER_P (get_kind (short_mapkind, kinds,
|
|
|
751 j)
|
|
|
752 & typemask))
|
|
|
753 break;
|
|
|
754 else if ((uintptr_t) hostaddrs[j] < k->host_start
|
|
|
755 || ((uintptr_t) hostaddrs[j] + sizeof (void *)
|
|
|
756 > k->host_end))
|
|
|
757 break;
|
|
|
758 else
|
|
|
759 {
|
|
|
760 tgt->list[j].key = k;
|
|
|
761 tgt->list[j].copy_from = false;
|
|
|
762 tgt->list[j].always_copy_from = false;
|
|
|
763 if (k->refcount != REFCOUNT_INFINITY)
|
|
|
764 k->refcount++;
|
|
|
765 gomp_map_pointer (tgt,
|
|
|
766 (uintptr_t) *(void **) hostaddrs[j],
|
|
|
767 k->tgt_offset
|
|
|
768 + ((uintptr_t) hostaddrs[j]
|
|
|
769 - k->host_start),
|
|
|
770 sizes[j]);
|
|
|
771 i++;
|
|
|
772 }
|
|
|
773 break;
|
|
|
774 case GOMP_MAP_FORCE_PRESENT:
|
|
|
775 {
|
|
|
776 /* We already looked up the memory region above and it
|
|
|
777 was missing. */
|
|
|
778 size_t size = k->host_end - k->host_start;
|
|
|
779 gomp_mutex_unlock (&devicep->lock);
|
|
|
780 #ifdef HAVE_INTTYPES_H
|
|
|
781 gomp_fatal ("present clause: !acc_is_present (%p, "
|
|
|
782 "%"PRIu64" (0x%"PRIx64"))",
|
|
|
783 (void *) k->host_start,
|
|
|
784 (uint64_t) size, (uint64_t) size);
|
|
|
785 #else
|
|
|
786 gomp_fatal ("present clause: !acc_is_present (%p, "
|
|
|
787 "%lu (0x%lx))", (void *) k->host_start,
|
|
|
788 (unsigned long) size, (unsigned long) size);
|
|
|
789 #endif
|
|
|
790 }
|
|
|
791 break;
|
|
|
792 case GOMP_MAP_FORCE_DEVICEPTR:
|
|
|
793 assert (k->host_end - k->host_start == sizeof (void *));
|
|
|
794 gomp_copy_host2dev (devicep,
|
|
|
795 (void *) (tgt->tgt_start
|
|
|
796 + k->tgt_offset),
|
|
|
797 (void *) k->host_start,
|
|
|
798 sizeof (void *));
|
|
|
799 break;
|
|
|
800 default:
|
|
|
801 gomp_mutex_unlock (&devicep->lock);
|
|
|
802 gomp_fatal ("%s: unhandled kind 0x%.2x", __FUNCTION__,
|
|
|
803 kind);
|
|
|
804 }
|
|
|
805
|
|
|
806 if (k->link_key)
|
|
|
807 {
|
|
|
808 /* Set link pointer on target to the device address of the
|
|
|
809 mapped object. */
|
|
|
810 void *tgt_addr = (void *) (tgt->tgt_start + k->tgt_offset);
|
|
|
811 devicep->host2dev_func (devicep->target_id,
|
|
|
812 (void *) n->tgt_offset,
|
|
|
813 &tgt_addr, sizeof (void *));
|
|
|
814 }
|
|
|
815 array++;
|
|
|
816 }
|
|
|
817 }
|
|
|
818 }
|
|
|
819
|
|
|
820 if (pragma_kind == GOMP_MAP_VARS_TARGET)
|
|
|
821 {
|
|
|
822 for (i = 0; i < mapnum; i++)
|
|
|
823 {
|
|
|
824 cur_node.tgt_offset = gomp_map_val (tgt, hostaddrs, i);
|
|
|
825 /* FIXME: see above FIXME comment. */
|
|
|
826 gomp_copy_host2dev (devicep,
|
|
|
827 (void *) (tgt->tgt_start + i * sizeof (void *)),
|
|
|
828 (void *) &cur_node.tgt_offset, sizeof (void *));
|
|
|
829 }
|
|
|
830 }
|
|
|
831
|
|
|
832 /* If the variable from "omp target enter data" map-list was already mapped,
|
|
|
833 tgt is not needed. Otherwise tgt will be freed by gomp_unmap_vars or
|
|
|
834 gomp_exit_data. */
|
|
|
835 if (pragma_kind == GOMP_MAP_VARS_ENTER_DATA && tgt->refcount == 0)
|
|
|
836 {
|
|
|
837 free (tgt);
|
|
|
838 tgt = NULL;
|
|
|
839 }
|
|
|
840
|
|
|
841 gomp_mutex_unlock (&devicep->lock);
|
|
|
842 return tgt;
|
|
|
843 }
|
|
|
844
|
|
|
845 static void
|
|
|
846 gomp_unmap_tgt (struct target_mem_desc *tgt)
|
|
|
847 {
|
|
|
848 /* Deallocate on target the tgt->tgt_start .. tgt->tgt_end region. */
|
|
|
849 if (tgt->tgt_end)
|
|
|
850 gomp_free_device_memory (tgt->device_descr, tgt->to_free);
|
|
|
851
|
|
|
852 free (tgt->array);
|
|
|
853 free (tgt);
|
|
|
854 }
|
|
|
855
|
|
|
856 /* Unmap variables described by TGT. If DO_COPYFROM is true, copy relevant
|
|
|
857 variables back from device to host: if it is false, it is assumed that this
|
|
|
858 has been done already. */
|
|
|
859
|
|
|
860 attribute_hidden void
|
|
|
861 gomp_unmap_vars (struct target_mem_desc *tgt, bool do_copyfrom)
|
|
|
862 {
|
|
|
863 struct gomp_device_descr *devicep = tgt->device_descr;
|
|
|
864
|
|
|
865 if (tgt->list_count == 0)
|
|
|
866 {
|
|
|
867 free (tgt);
|
|
|
868 return;
|
|
|
869 }
|
|
|
870
|
|
|
871 gomp_mutex_lock (&devicep->lock);
|
|
|
872 if (devicep->state == GOMP_DEVICE_FINALIZED)
|
|
|
873 {
|
|
|
874 gomp_mutex_unlock (&devicep->lock);
|
|
|
875 free (tgt->array);
|
|
|
876 free (tgt);
|
|
|
877 return;
|
|
|
878 }
|
|
|
879
|
|
|
880 size_t i;
|
|
|
881 for (i = 0; i < tgt->list_count; i++)
|
|
|
882 {
|
|
|
883 splay_tree_key k = tgt->list[i].key;
|
|
|
884 if (k == NULL)
|
|
|
885 continue;
|
|
|
886
|
|
|
887 bool do_unmap = false;
|
|
|
888 if (k->refcount > 1 && k->refcount != REFCOUNT_INFINITY)
|
|
|
889 k->refcount--;
|
|
|
890 else if (k->refcount == 1)
|
|
|
891 {
|
|
|
892 k->refcount--;
|
|
|
893 do_unmap = true;
|
|
|
894 }
|
|
|
895
|
|
|
896 if ((do_unmap && do_copyfrom && tgt->list[i].copy_from)
|
|
|
897 || tgt->list[i].always_copy_from)
|
|
|
898 gomp_copy_dev2host (devicep,
|
|
|
899 (void *) (k->host_start + tgt->list[i].offset),
|
|
|
900 (void *) (k->tgt->tgt_start + k->tgt_offset
|
|
|
901 + tgt->list[i].offset),
|
|
|
902 tgt->list[i].length);
|
|
|
903 if (do_unmap)
|
|
|
904 {
|
|
|
905 splay_tree_remove (&devicep->mem_map, k);
|
|
|
906 if (k->link_key)
|
|
|
907 splay_tree_insert (&devicep->mem_map,
|
|
|
908 (splay_tree_node) k->link_key);
|
|
|
909 if (k->tgt->refcount > 1)
|
|
|
910 k->tgt->refcount--;
|
|
|
911 else
|
|
|
912 gomp_unmap_tgt (k->tgt);
|
|
|
913 }
|
|
|
914 }
|
|
|
915
|
|
|
916 if (tgt->refcount > 1)
|
|
|
917 tgt->refcount--;
|
|
|
918 else
|
|
|
919 gomp_unmap_tgt (tgt);
|
|
|
920
|
|
|
921 gomp_mutex_unlock (&devicep->lock);
|
|
|
922 }
|
|
|
923
|
|
|
924 static void
|
|
|
925 gomp_update (struct gomp_device_descr *devicep, size_t mapnum, void **hostaddrs,
|
|
|
926 size_t *sizes, void *kinds, bool short_mapkind)
|
|
|
927 {
|
|
|
928 size_t i;
|
|
|
929 struct splay_tree_key_s cur_node;
|
|
|
930 const int typemask = short_mapkind ? 0xff : 0x7;
|
|
|
931
|
|
|
932 if (!devicep)
|
|
|
933 return;
|
|
|
934
|
|
|
935 if (mapnum == 0)
|
|
|
936 return;
|
|
|
937
|
|
|
938 gomp_mutex_lock (&devicep->lock);
|
|
|
939 if (devicep->state == GOMP_DEVICE_FINALIZED)
|
|
|
940 {
|
|
|
941 gomp_mutex_unlock (&devicep->lock);
|
|
|
942 return;
|
|
|
943 }
|
|
|
944
|
|
|
945 for (i = 0; i < mapnum; i++)
|
|
|
946 if (sizes[i])
|
|
|
947 {
|
|
|
948 cur_node.host_start = (uintptr_t) hostaddrs[i];
|
|
|
949 cur_node.host_end = cur_node.host_start + sizes[i];
|
|
|
950 splay_tree_key n = splay_tree_lookup (&devicep->mem_map, &cur_node);
|
|
|
951 if (n)
|
|
|
952 {
|
|
|
953 int kind = get_kind (short_mapkind, kinds, i);
|
|
|
954 if (n->host_start > cur_node.host_start
|
|
|
955 || n->host_end < cur_node.host_end)
|
|
|
956 {
|
|
|
957 gomp_mutex_unlock (&devicep->lock);
|
|
|
958 gomp_fatal ("Trying to update [%p..%p) object when "
|
|
|
959 "only [%p..%p) is mapped",
|
|
|
960 (void *) cur_node.host_start,
|
|
|
961 (void *) cur_node.host_end,
|
|
|
962 (void *) n->host_start,
|
|
|
963 (void *) n->host_end);
|
|
|
964 }
|
|
|
965
|
|
|
966
|
|
|
967 void *hostaddr = (void *) cur_node.host_start;
|
|
|
968 void *devaddr = (void *) (n->tgt->tgt_start + n->tgt_offset
|
|
|
969 + cur_node.host_start - n->host_start);
|
|
|
970 size_t size = cur_node.host_end - cur_node.host_start;
|
|
|
971
|
|
|
972 if (GOMP_MAP_COPY_TO_P (kind & typemask))
|
|
|
973 gomp_copy_host2dev (devicep, devaddr, hostaddr, size);
|
|
|
974 if (GOMP_MAP_COPY_FROM_P (kind & typemask))
|
|
|
975 gomp_copy_dev2host (devicep, hostaddr, devaddr, size);
|
|
|
976 }
|
|
|
977 }
|
|
|
978 gomp_mutex_unlock (&devicep->lock);
|
|
|
979 }
|
|
|
980
|
|
|
981 /* Load image pointed by TARGET_DATA to the device, specified by DEVICEP.
|
|
|
982 And insert to splay tree the mapping between addresses from HOST_TABLE and
|
|
|
983 from loaded target image. We rely in the host and device compiler
|
|
|
984 emitting variable and functions in the same order. */
|
|
|
985
|
|
|
986 static void
|
|
|
987 gomp_load_image_to_device (struct gomp_device_descr *devicep, unsigned version,
|
|
|
988 const void *host_table, const void *target_data,
|
|
|
989 bool is_register_lock)
|
|
|
990 {
|
|
|
991 void **host_func_table = ((void ***) host_table)[0];
|
|
|
992 void **host_funcs_end = ((void ***) host_table)[1];
|
|
|
993 void **host_var_table = ((void ***) host_table)[2];
|
|
|
994 void **host_vars_end = ((void ***) host_table)[3];
|
|
|
995
|
|
|
996 /* The func table contains only addresses, the var table contains addresses
|
|
|
997 and corresponding sizes. */
|
|
|
998 int num_funcs = host_funcs_end - host_func_table;
|
|
|
999 int num_vars = (host_vars_end - host_var_table) / 2;
|
|
|
1000
|
|
|
1001 /* Load image to device and get target addresses for the image. */
|
|
|
1002 struct addr_pair *target_table = NULL;
|
|
|
1003 int i, num_target_entries;
|
|
|
1004
|
|
|
1005 num_target_entries
|
|
|
1006 = devicep->load_image_func (devicep->target_id, version,
|
|
|
1007 target_data, &target_table);
|
|
|
1008
|
|
|
1009 if (num_target_entries != num_funcs + num_vars)
|
|
|
1010 {
|
|
|
1011 gomp_mutex_unlock (&devicep->lock);
|
|
|
1012 if (is_register_lock)
|
|
|
1013 gomp_mutex_unlock (®ister_lock);
|
|
|
1014 gomp_fatal ("Cannot map target functions or variables"
|
|
|
1015 " (expected %u, have %u)", num_funcs + num_vars,
|
|
|
1016 num_target_entries);
|
|
|
1017 }
|
|
|
1018
|
|
|
1019 /* Insert host-target address mapping into splay tree. */
|
|
|
1020 struct target_mem_desc *tgt = gomp_malloc (sizeof (*tgt));
|
|
|
1021 tgt->array = gomp_malloc ((num_funcs + num_vars) * sizeof (*tgt->array));
|
|
|
1022 tgt->refcount = REFCOUNT_INFINITY;
|
|
|
1023 tgt->tgt_start = 0;
|
|
|
1024 tgt->tgt_end = 0;
|
|
|
1025 tgt->to_free = NULL;
|
|
|
1026 tgt->prev = NULL;
|
|
|
1027 tgt->list_count = 0;
|
|
|
1028 tgt->device_descr = devicep;
|
|
|
1029 splay_tree_node array = tgt->array;
|
|
|
1030
|
|
|
1031 for (i = 0; i < num_funcs; i++)
|
|
|
1032 {
|
|
|
1033 splay_tree_key k = &array->key;
|
|
|
1034 k->host_start = (uintptr_t) host_func_table[i];
|
|
|
1035 k->host_end = k->host_start + 1;
|
|
|
1036 k->tgt = tgt;
|
|
|
1037 k->tgt_offset = target_table[i].start;
|
|
|
1038 k->refcount = REFCOUNT_INFINITY;
|
|
|
1039 k->link_key = NULL;
|
|
|
1040 array->left = NULL;
|
|
|
1041 array->right = NULL;
|
|
|
1042 splay_tree_insert (&devicep->mem_map, array);
|
|
|
1043 array++;
|
|
|
1044 }
|
|
|
1045
|
|
|
1046 /* Most significant bit of the size in host and target tables marks
|
|
|
1047 "omp declare target link" variables. */
|
|
|
1048 const uintptr_t link_bit = 1ULL << (sizeof (uintptr_t) * __CHAR_BIT__ - 1);
|
|
|
1049 const uintptr_t size_mask = ~link_bit;
|
|
|
1050
|
|
|
1051 for (i = 0; i < num_vars; i++)
|
|
|
1052 {
|
|
|
1053 struct addr_pair *target_var = &target_table[num_funcs + i];
|
|
|
1054 uintptr_t target_size = target_var->end - target_var->start;
|
|
|
1055
|
|
|
1056 if ((uintptr_t) host_var_table[i * 2 + 1] != target_size)
|
|
|
1057 {
|
|
|
1058 gomp_mutex_unlock (&devicep->lock);
|
|
|
1059 if (is_register_lock)
|
|
|
1060 gomp_mutex_unlock (®ister_lock);
|
|
|
1061 gomp_fatal ("Cannot map target variables (size mismatch)");
|
|
|
1062 }
|
|
|
1063
|
|
|
1064 splay_tree_key k = &array->key;
|
|
|
1065 k->host_start = (uintptr_t) host_var_table[i * 2];
|
|
|
1066 k->host_end
|
|
|
1067 = k->host_start + (size_mask & (uintptr_t) host_var_table[i * 2 + 1]);
|
|
|
1068 k->tgt = tgt;
|
|
|
1069 k->tgt_offset = target_var->start;
|
|
|
1070 k->refcount = target_size & link_bit ? REFCOUNT_LINK : REFCOUNT_INFINITY;
|
|
|
1071 k->link_key = NULL;
|
|
|
1072 array->left = NULL;
|
|
|
1073 array->right = NULL;
|
|
|
1074 splay_tree_insert (&devicep->mem_map, array);
|
|
|
1075 array++;
|
|
|
1076 }
|
|
|
1077
|
|
|
1078 free (target_table);
|
|
|
1079 }
|
|
|
1080
|
|
|
1081 /* Unload the mappings described by target_data from device DEVICE_P.
|
|
|
1082 The device must be locked. */
|
|
|
1083
|
|
|
1084 static void
|
|
|
1085 gomp_unload_image_from_device (struct gomp_device_descr *devicep,
|
|
|
1086 unsigned version,
|
|
|
1087 const void *host_table, const void *target_data)
|
|
|
1088 {
|
|
|
1089 void **host_func_table = ((void ***) host_table)[0];
|
|
|
1090 void **host_funcs_end = ((void ***) host_table)[1];
|
|
|
1091 void **host_var_table = ((void ***) host_table)[2];
|
|
|
1092 void **host_vars_end = ((void ***) host_table)[3];
|
|
|
1093
|
|
|
1094 /* The func table contains only addresses, the var table contains addresses
|
|
|
1095 and corresponding sizes. */
|
|
|
1096 int num_funcs = host_funcs_end - host_func_table;
|
|
|
1097 int num_vars = (host_vars_end - host_var_table) / 2;
|
|
|
1098
|
|
|
1099 struct splay_tree_key_s k;
|
|
|
1100 splay_tree_key node = NULL;
|
|
|
1101
|
|
|
1102 /* Find mapping at start of node array */
|
|
|
1103 if (num_funcs || num_vars)
|
|
|
1104 {
|
|
|
1105 k.host_start = (num_funcs ? (uintptr_t) host_func_table[0]
|
|
|
1106 : (uintptr_t) host_var_table[0]);
|
|
|
1107 k.host_end = k.host_start + 1;
|
|
|
1108 node = splay_tree_lookup (&devicep->mem_map, &k);
|
|
|
1109 }
|
|
|
1110
|
|
|
1111 if (!devicep->unload_image_func (devicep->target_id, version, target_data))
|
|
|
1112 {
|
|
|
1113 gomp_mutex_unlock (&devicep->lock);
|
|
|
1114 gomp_fatal ("image unload fail");
|
|
|
1115 }
|
|
|
1116
|
|
|
1117 /* Remove mappings from splay tree. */
|
|
|
1118 int i;
|
|
|
1119 for (i = 0; i < num_funcs; i++)
|
|
|
1120 {
|
|
|
1121 k.host_start = (uintptr_t) host_func_table[i];
|
|
|
1122 k.host_end = k.host_start + 1;
|
|
|
1123 splay_tree_remove (&devicep->mem_map, &k);
|
|
|
1124 }
|
|
|
1125
|
|
|
1126 /* Most significant bit of the size in host and target tables marks
|
|
|
1127 "omp declare target link" variables. */
|
|
|
1128 const uintptr_t link_bit = 1ULL << (sizeof (uintptr_t) * __CHAR_BIT__ - 1);
|
|
|
1129 const uintptr_t size_mask = ~link_bit;
|
|
|
1130 bool is_tgt_unmapped = false;
|
|
|
1131
|
|
|
1132 for (i = 0; i < num_vars; i++)
|
|
|
1133 {
|
|
|
1134 k.host_start = (uintptr_t) host_var_table[i * 2];
|
|
|
1135 k.host_end
|
|
|
1136 = k.host_start + (size_mask & (uintptr_t) host_var_table[i * 2 + 1]);
|
|
|
1137
|
|
|
1138 if (!(link_bit & (uintptr_t) host_var_table[i * 2 + 1]))
|
|
|
1139 splay_tree_remove (&devicep->mem_map, &k);
|
|
|
1140 else
|
|
|
1141 {
|
|
|
1142 splay_tree_key n = splay_tree_lookup (&devicep->mem_map, &k);
|
|
|
1143 splay_tree_remove (&devicep->mem_map, n);
|
|
|
1144 if (n->link_key)
|
|
|
1145 {
|
|
|
1146 if (n->tgt->refcount > 1)
|
|
|
1147 n->tgt->refcount--;
|
|
|
1148 else
|
|
|
1149 {
|
|
|
1150 is_tgt_unmapped = true;
|
|
|
1151 gomp_unmap_tgt (n->tgt);
|
|
|
1152 }
|
|
|
1153 }
|
|
|
1154 }
|
|
|
1155 }
|
|
|
1156
|
|
|
1157 if (node && !is_tgt_unmapped)
|
|
|
1158 {
|
|
|
1159 free (node->tgt);
|
|
|
1160 free (node);
|
|
|
1161 }
|
|
|
1162 }
|
|
|
1163
|
|
|
1164 /* This function should be called from every offload image while loading.
|
|
|
1165 It gets the descriptor of the host func and var tables HOST_TABLE, TYPE of
|
|
|
1166 the target, and TARGET_DATA needed by target plugin. */
|
|
|
1167
|
|
|
1168 void
|
|
|
1169 GOMP_offload_register_ver (unsigned version, const void *host_table,
|
|
|
1170 int target_type, const void *target_data)
|
|
|
1171 {
|
|
|
1172 int i;
|
|
|
1173
|
|
|
1174 if (GOMP_VERSION_LIB (version) > GOMP_VERSION)
|
|
|
1175 gomp_fatal ("Library too old for offload (version %u < %u)",
|
|
|
1176 GOMP_VERSION, GOMP_VERSION_LIB (version));
|
|
|
1177
|
|
|
1178 gomp_mutex_lock (®ister_lock);
|
|
|
1179
|
|
|
1180 /* Load image to all initialized devices. */
|
|
|
1181 for (i = 0; i < num_devices; i++)
|
|
|
1182 {
|
|
|
1183 struct gomp_device_descr *devicep = &devices[i];
|
|
|
1184 gomp_mutex_lock (&devicep->lock);
|
|
|
1185 if (devicep->type == target_type
|
|
|
1186 && devicep->state == GOMP_DEVICE_INITIALIZED)
|
|
|
1187 gomp_load_image_to_device (devicep, version,
|
|
|
1188 host_table, target_data, true);
|
|
|
1189 gomp_mutex_unlock (&devicep->lock);
|
|
|
1190 }
|
|
|
1191
|
|
|
1192 /* Insert image to array of pending images. */
|
|
|
1193 offload_images
|
|
|
1194 = gomp_realloc_unlock (offload_images,
|
|
|
1195 (num_offload_images + 1)
|
|
|
1196 * sizeof (struct offload_image_descr));
|
|
|
1197 offload_images[num_offload_images].version = version;
|
|
|
1198 offload_images[num_offload_images].type = target_type;
|
|
|
1199 offload_images[num_offload_images].host_table = host_table;
|
|
|
1200 offload_images[num_offload_images].target_data = target_data;
|
|
|
1201
|
|
|
1202 num_offload_images++;
|
|
|
1203 gomp_mutex_unlock (®ister_lock);
|
|
|
1204 }
|
|
|
1205
|
|
|
1206 void
|
|
|
1207 GOMP_offload_register (const void *host_table, int target_type,
|
|
|
1208 const void *target_data)
|
|
|
1209 {
|
|
|
1210 GOMP_offload_register_ver (0, host_table, target_type, target_data);
|
|
|
1211 }
|
|
|
1212
|
|
|
1213 /* This function should be called from every offload image while unloading.
|
|
|
1214 It gets the descriptor of the host func and var tables HOST_TABLE, TYPE of
|
|
|
1215 the target, and TARGET_DATA needed by target plugin. */
|
|
|
1216
|
|
|
1217 void
|
|
|
1218 GOMP_offload_unregister_ver (unsigned version, const void *host_table,
|
|
|
1219 int target_type, const void *target_data)
|
|
|
1220 {
|
|
|
1221 int i;
|
|
|
1222
|
|
|
1223 gomp_mutex_lock (®ister_lock);
|
|
|
1224
|
|
|
1225 /* Unload image from all initialized devices. */
|
|
|
1226 for (i = 0; i < num_devices; i++)
|
|
|
1227 {
|
|
|
1228 struct gomp_device_descr *devicep = &devices[i];
|
|
|
1229 gomp_mutex_lock (&devicep->lock);
|
|
|
1230 if (devicep->type == target_type
|
|
|
1231 && devicep->state == GOMP_DEVICE_INITIALIZED)
|
|
|
1232 gomp_unload_image_from_device (devicep, version,
|
|
|
1233 host_table, target_data);
|
|
|
1234 gomp_mutex_unlock (&devicep->lock);
|
|
|
1235 }
|
|
|
1236
|
|
|
1237 /* Remove image from array of pending images. */
|
|
|
1238 for (i = 0; i < num_offload_images; i++)
|
|
|
1239 if (offload_images[i].target_data == target_data)
|
|
|
1240 {
|
|
|
1241 offload_images[i] = offload_images[--num_offload_images];
|
|
|
1242 break;
|
|
|
1243 }
|
|
|
1244
|
|
|
1245 gomp_mutex_unlock (®ister_lock);
|
|
|
1246 }
|
|
|
1247
|
|
|
1248 void
|
|
|
1249 GOMP_offload_unregister (const void *host_table, int target_type,
|
|
|
1250 const void *target_data)
|
|
|
1251 {
|
|
|
1252 GOMP_offload_unregister_ver (0, host_table, target_type, target_data);
|
|
|
1253 }
|
|
|
1254
|
|
|
1255 /* This function initializes the target device, specified by DEVICEP. DEVICEP
|
|
|
1256 must be locked on entry, and remains locked on return. */
|
|
|
1257
|
|
|
1258 attribute_hidden void
|
|
|
1259 gomp_init_device (struct gomp_device_descr *devicep)
|
|
|
1260 {
|
|
|
1261 int i;
|
|
|
1262 if (!devicep->init_device_func (devicep->target_id))
|
|
|
1263 {
|
|
|
1264 gomp_mutex_unlock (&devicep->lock);
|
|
|
1265 gomp_fatal ("device initialization failed");
|
|
|
1266 }
|
|
|
1267
|
|
|
1268 /* Load to device all images registered by the moment. */
|
|
|
1269 for (i = 0; i < num_offload_images; i++)
|
|
|
1270 {
|
|
|
1271 struct offload_image_descr *image = &offload_images[i];
|
|
|
1272 if (image->type == devicep->type)
|
|
|
1273 gomp_load_image_to_device (devicep, image->version,
|
|
|
1274 image->host_table, image->target_data,
|
|
|
1275 false);
|
|
|
1276 }
|
|
|
1277
|
|
|
1278 devicep->state = GOMP_DEVICE_INITIALIZED;
|
|
|
1279 }
|
|
|
1280
|
|
|
1281 attribute_hidden void
|
|
|
1282 gomp_unload_device (struct gomp_device_descr *devicep)
|
|
|
1283 {
|
|
|
1284 if (devicep->state == GOMP_DEVICE_INITIALIZED)
|
|
|
1285 {
|
|
|
1286 unsigned i;
|
|
|
1287
|
|
|
1288 /* Unload from device all images registered at the moment. */
|
|
|
1289 for (i = 0; i < num_offload_images; i++)
|
|
|
1290 {
|
|
|
1291 struct offload_image_descr *image = &offload_images[i];
|
|
|
1292 if (image->type == devicep->type)
|
|
|
1293 gomp_unload_image_from_device (devicep, image->version,
|
|
|
1294 image->host_table,
|
|
|
1295 image->target_data);
|
|
|
1296 }
|
|
|
1297 }
|
|
|
1298 }
|
|
|
1299
|
|
|
1300 /* Free address mapping tables. MM must be locked on entry, and remains locked
|
|
|
1301 on return. */
|
|
|
1302
|
|
|
1303 attribute_hidden void
|
|
|
1304 gomp_free_memmap (struct splay_tree_s *mem_map)
|
|
|
1305 {
|
|
|
1306 while (mem_map->root)
|
|
|
1307 {
|
|
|
1308 struct target_mem_desc *tgt = mem_map->root->key.tgt;
|
|
|
1309
|
|
|
1310 splay_tree_remove (mem_map, &mem_map->root->key);
|
|
|
1311 free (tgt->array);
|
|
|
1312 free (tgt);
|
|
|
1313 }
|
|
|
1314 }
|
|
|
1315
|
|
|
1316 /* Host fallback for GOMP_target{,_ext} routines. */
|
|
|
1317
|
|
|
1318 static void
|
|
|
1319 gomp_target_fallback (void (*fn) (void *), void **hostaddrs)
|
|
|
1320 {
|
|
|
1321 struct gomp_thread old_thr, *thr = gomp_thread ();
|
|
|
1322 old_thr = *thr;
|
|
|
1323 memset (thr, '\0', sizeof (*thr));
|
|
|
1324 if (gomp_places_list)
|
|
|
1325 {
|
|
|
1326 thr->place = old_thr.place;
|
|
|
1327 thr->ts.place_partition_len = gomp_places_list_len;
|
|
|
1328 }
|
|
|
1329 fn (hostaddrs);
|
|
|
1330 gomp_free_thread (thr);
|
|
|
1331 *thr = old_thr;
|
|
|
1332 }
|
|
|
1333
|
|
|
1334 /* Calculate alignment and size requirements of a private copy of data shared
|
|
|
1335 as GOMP_MAP_FIRSTPRIVATE and store them to TGT_ALIGN and TGT_SIZE. */
|
|
|
1336
|
|
|
1337 static inline void
|
|
|
1338 calculate_firstprivate_requirements (size_t mapnum, size_t *sizes,
|
|
|
1339 unsigned short *kinds, size_t *tgt_align,
|
|
|
1340 size_t *tgt_size)
|
|
|
1341 {
|
|
|
1342 size_t i;
|
|
|
1343 for (i = 0; i < mapnum; i++)
|
|
|
1344 if ((kinds[i] & 0xff) == GOMP_MAP_FIRSTPRIVATE)
|
|
|
1345 {
|
|
|
1346 size_t align = (size_t) 1 << (kinds[i] >> 8);
|
|
|
1347 if (*tgt_align < align)
|
|
|
1348 *tgt_align = align;
|
|
|
1349 *tgt_size = (*tgt_size + align - 1) & ~(align - 1);
|
|
|
1350 *tgt_size += sizes[i];
|
|
|
1351 }
|
|
|
1352 }
|
|
|
1353
|
|
|
1354 /* Copy data shared as GOMP_MAP_FIRSTPRIVATE to DST. */
|
|
|
1355
|
|
|
1356 static inline void
|
|
|
1357 copy_firstprivate_data (char *tgt, size_t mapnum, void **hostaddrs,
|
|
|
1358 size_t *sizes, unsigned short *kinds, size_t tgt_align,
|
|
|
1359 size_t tgt_size)
|
|
|
1360 {
|
|
|
1361 uintptr_t al = (uintptr_t) tgt & (tgt_align - 1);
|
|
|
1362 if (al)
|
|
|
1363 tgt += tgt_align - al;
|
|
|
1364 tgt_size = 0;
|
|
|
1365 size_t i;
|
|
|
1366 for (i = 0; i < mapnum; i++)
|
|
|
1367 if ((kinds[i] & 0xff) == GOMP_MAP_FIRSTPRIVATE)
|
|
|
1368 {
|
|
|
1369 size_t align = (size_t) 1 << (kinds[i] >> 8);
|
|
|
1370 tgt_size = (tgt_size + align - 1) & ~(align - 1);
|
|
|
1371 memcpy (tgt + tgt_size, hostaddrs[i], sizes[i]);
|
|
|
1372 hostaddrs[i] = tgt + tgt_size;
|
|
|
1373 tgt_size = tgt_size + sizes[i];
|
|
|
1374 }
|
|
|
1375 }
|
|
|
1376
|
|
|
1377 /* Helper function of GOMP_target{,_ext} routines. */
|
|
|
1378
|
|
|
1379 static void *
|
|
|
1380 gomp_get_target_fn_addr (struct gomp_device_descr *devicep,
|
|
|
1381 void (*host_fn) (void *))
|
|
|
1382 {
|
|
|
1383 if (devicep->capabilities & GOMP_OFFLOAD_CAP_NATIVE_EXEC)
|
|
|
1384 return (void *) host_fn;
|
|
|
1385 else
|
|
|
1386 {
|
|
|
1387 gomp_mutex_lock (&devicep->lock);
|
|
|
1388 if (devicep->state == GOMP_DEVICE_FINALIZED)
|
|
|
1389 {
|
|
|
1390 gomp_mutex_unlock (&devicep->lock);
|
|
|
1391 return NULL;
|
|
|
1392 }
|
|
|
1393
|
|
|
1394 struct splay_tree_key_s k;
|
|
|
1395 k.host_start = (uintptr_t) host_fn;
|
|
|
1396 k.host_end = k.host_start + 1;
|
|
|
1397 splay_tree_key tgt_fn = splay_tree_lookup (&devicep->mem_map, &k);
|
|
|
1398 gomp_mutex_unlock (&devicep->lock);
|
|
|
1399 if (tgt_fn == NULL)
|
|
|
1400 return NULL;
|
|
|
1401
|
|
|
1402 return (void *) tgt_fn->tgt_offset;
|
|
|
1403 }
|
|
|
1404 }
|
|
|
1405
|
|
|
1406 /* Called when encountering a target directive. If DEVICE
|
|
|
1407 is GOMP_DEVICE_ICV, it means use device-var ICV. If it is
|
|
|
1408 GOMP_DEVICE_HOST_FALLBACK (or any value
|
|
|
1409 larger than last available hw device), use host fallback.
|
|
|
1410 FN is address of host code, UNUSED is part of the current ABI, but
|
|
|
1411 we're not actually using it. HOSTADDRS, SIZES and KINDS are arrays
|
|
|
1412 with MAPNUM entries, with addresses of the host objects,
|
|
|
1413 sizes of the host objects (resp. for pointer kind pointer bias
|
|
|
1414 and assumed sizeof (void *) size) and kinds. */
|
|
|
1415
|
|
|
1416 void
|
|
|
1417 GOMP_target (int device, void (*fn) (void *), const void *unused,
|
|
|
1418 size_t mapnum, void **hostaddrs, size_t *sizes,
|
|
|
1419 unsigned char *kinds)
|
|
|
1420 {
|
|
|
1421 struct gomp_device_descr *devicep = resolve_device (device);
|
|
|
1422
|
|
|
1423 void *fn_addr;
|
|
|
1424 if (devicep == NULL
|
|
|
1425 || !(devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400)
|
|
|
1426 /* All shared memory devices should use the GOMP_target_ext function. */
|
|
|
1427 || devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM
|
|
|
1428 || !(fn_addr = gomp_get_target_fn_addr (devicep, fn)))
|
|
|
1429 return gomp_target_fallback (fn, hostaddrs);
|
|
|
1430
|
|
|
1431 struct target_mem_desc *tgt_vars
|
|
|
1432 = gomp_map_vars (devicep, mapnum, hostaddrs, NULL, sizes, kinds, false,
|
|
|
1433 GOMP_MAP_VARS_TARGET);
|
|
|
1434 devicep->run_func (devicep->target_id, fn_addr, (void *) tgt_vars->tgt_start,
|
|
|
1435 NULL);
|
|
|
1436 gomp_unmap_vars (tgt_vars, true);
|
|
|
1437 }
|
|
|
1438
|
|
|
1439 /* Like GOMP_target, but KINDS is 16-bit, UNUSED is no longer present,
|
|
|
1440 and several arguments have been added:
|
|
|
1441 FLAGS is a bitmask, see GOMP_TARGET_FLAG_* in gomp-constants.h.
|
|
|
1442 DEPEND is array of dependencies, see GOMP_task for details.
|
|
|
1443
|
|
|
1444 ARGS is a pointer to an array consisting of a variable number of both
|
|
|
1445 device-independent and device-specific arguments, which can take one two
|
|
|
1446 elements where the first specifies for which device it is intended, the type
|
|
|
1447 and optionally also the value. If the value is not present in the first
|
|
|
1448 one, the whole second element the actual value. The last element of the
|
|
|
1449 array is a single NULL. Among the device independent can be for example
|
|
|
1450 NUM_TEAMS and THREAD_LIMIT.
|
|
|
1451
|
|
|
1452 NUM_TEAMS is positive if GOMP_teams will be called in the body with
|
|
|
1453 that value, or 1 if teams construct is not present, or 0, if
|
|
|
1454 teams construct does not have num_teams clause and so the choice is
|
|
|
1455 implementation defined, and -1 if it can't be determined on the host
|
|
|
1456 what value will GOMP_teams have on the device.
|
|
|
1457 THREAD_LIMIT similarly is positive if GOMP_teams will be called in the
|
|
|
1458 body with that value, or 0, if teams construct does not have thread_limit
|
|
|
1459 clause or the teams construct is not present, or -1 if it can't be
|
|
|
1460 determined on the host what value will GOMP_teams have on the device. */
|
|
|
1461
|
|
|
1462 void
|
|
|
1463 GOMP_target_ext (int device, void (*fn) (void *), size_t mapnum,
|
|
|
1464 void **hostaddrs, size_t *sizes, unsigned short *kinds,
|
|
|
1465 unsigned int flags, void **depend, void **args)
|
|
|
1466 {
|
|
|
1467 struct gomp_device_descr *devicep = resolve_device (device);
|
|
|
1468 size_t tgt_align = 0, tgt_size = 0;
|
|
|
1469 bool fpc_done = false;
|
|
|
1470
|
|
|
1471 if (flags & GOMP_TARGET_FLAG_NOWAIT)
|
|
|
1472 {
|
|
|
1473 struct gomp_thread *thr = gomp_thread ();
|
|
|
1474 /* Create a team if we don't have any around, as nowait
|
|
|
1475 target tasks make sense to run asynchronously even when
|
|
|
1476 outside of any parallel. */
|
|
|
1477 if (__builtin_expect (thr->ts.team == NULL, 0))
|
|
|
1478 {
|
|
|
1479 struct gomp_team *team = gomp_new_team (1);
|
|
|
1480 struct gomp_task *task = thr->task;
|
|
|
1481 struct gomp_task_icv *icv = task ? &task->icv : &gomp_global_icv;
|
|
|
1482 team->prev_ts = thr->ts;
|
|
|
1483 thr->ts.team = team;
|
|
|
1484 thr->ts.team_id = 0;
|
|
|
1485 thr->ts.work_share = &team->work_shares[0];
|
|
|
1486 thr->ts.last_work_share = NULL;
|
|
|
1487 #ifdef HAVE_SYNC_BUILTINS
|
|
|
1488 thr->ts.single_count = 0;
|
|
|
1489 #endif
|
|
|
1490 thr->ts.static_trip = 0;
|
|
|
1491 thr->task = &team->implicit_task[0];
|
|
|
1492 gomp_init_task (thr->task, NULL, icv);
|
|
|
1493 if (task)
|
|
|
1494 {
|
|
|
1495 thr->task = task;
|
|
|
1496 gomp_end_task ();
|
|
|
1497 free (task);
|
|
|
1498 thr->task = &team->implicit_task[0];
|
|
|
1499 }
|
|
|
1500 else
|
|
|
1501 pthread_setspecific (gomp_thread_destructor, thr);
|
|
|
1502 }
|
|
|
1503 if (thr->ts.team
|
|
|
1504 && !thr->task->final_task)
|
|
|
1505 {
|
|
|
1506 gomp_create_target_task (devicep, fn, mapnum, hostaddrs,
|
|
|
1507 sizes, kinds, flags, depend, args,
|
|
|
1508 GOMP_TARGET_TASK_BEFORE_MAP);
|
|
|
1509 return;
|
|
|
1510 }
|
|
|
1511 }
|
|
|
1512
|
|
|
1513 /* If there are depend clauses, but nowait is not present
|
|
|
1514 (or we are in a final task), block the parent task until the
|
|
|
1515 dependencies are resolved and then just continue with the rest
|
|
|
1516 of the function as if it is a merged task. */
|
|
|
1517 if (depend != NULL)
|
|
|
1518 {
|
|
|
1519 struct gomp_thread *thr = gomp_thread ();
|
|
|
1520 if (thr->task && thr->task->depend_hash)
|
|
|
1521 {
|
|
|
1522 /* If we might need to wait, copy firstprivate now. */
|
|
|
1523 calculate_firstprivate_requirements (mapnum, sizes, kinds,
|
|
|
1524 &tgt_align, &tgt_size);
|
|
|
1525 if (tgt_align)
|
|
|
1526 {
|
|
|
1527 char *tgt = gomp_alloca (tgt_size + tgt_align - 1);
|
|
|
1528 copy_firstprivate_data (tgt, mapnum, hostaddrs, sizes, kinds,
|
|
|
1529 tgt_align, tgt_size);
|
|
|
1530 }
|
|
|
1531 fpc_done = true;
|
|
|
1532 gomp_task_maybe_wait_for_dependencies (depend);
|
|
|
1533 }
|
|
|
1534 }
|
|
|
1535
|
|
|
1536 void *fn_addr;
|
|
|
1537 if (devicep == NULL
|
|
|
1538 || !(devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400)
|
|
|
1539 || !(fn_addr = gomp_get_target_fn_addr (devicep, fn))
|
|
|
1540 || (devicep->can_run_func && !devicep->can_run_func (fn_addr)))
|
|
|
1541 {
|
|
|
1542 if (!fpc_done)
|
|
|
1543 {
|
|
|
1544 calculate_firstprivate_requirements (mapnum, sizes, kinds,
|
|
|
1545 &tgt_align, &tgt_size);
|
|
|
1546 if (tgt_align)
|
|
|
1547 {
|
|
|
1548 char *tgt = gomp_alloca (tgt_size + tgt_align - 1);
|
|
|
1549 copy_firstprivate_data (tgt, mapnum, hostaddrs, sizes, kinds,
|
|
|
1550 tgt_align, tgt_size);
|
|
|
1551 }
|
|
|
1552 }
|
|
|
1553 gomp_target_fallback (fn, hostaddrs);
|
|
|
1554 return;
|
|
|
1555 }
|
|
|
1556
|
|
|
1557 struct target_mem_desc *tgt_vars;
|
|
|
1558 if (devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
|
|
|
1559 {
|
|
|
1560 if (!fpc_done)
|
|
|
1561 {
|
|
|
1562 calculate_firstprivate_requirements (mapnum, sizes, kinds,
|
|
|
1563 &tgt_align, &tgt_size);
|
|
|
1564 if (tgt_align)
|
|
|
1565 {
|
|
|
1566 char *tgt = gomp_alloca (tgt_size + tgt_align - 1);
|
|
|
1567 copy_firstprivate_data (tgt, mapnum, hostaddrs, sizes, kinds,
|
|
|
1568 tgt_align, tgt_size);
|
|
|
1569 }
|
|
|
1570 }
|
|
|
1571 tgt_vars = NULL;
|
|
|
1572 }
|
|
|
1573 else
|
|
|
1574 tgt_vars = gomp_map_vars (devicep, mapnum, hostaddrs, NULL, sizes, kinds,
|
|
|
1575 true, GOMP_MAP_VARS_TARGET);
|
|
|
1576 devicep->run_func (devicep->target_id, fn_addr,
|
|
|
1577 tgt_vars ? (void *) tgt_vars->tgt_start : hostaddrs,
|
|
|
1578 args);
|
|
|
1579 if (tgt_vars)
|
|
|
1580 gomp_unmap_vars (tgt_vars, true);
|
|
|
1581 }
|
|
|
1582
|
|
|
1583 /* Host fallback for GOMP_target_data{,_ext} routines. */
|
|
|
1584
|
|
|
1585 static void
|
|
|
1586 gomp_target_data_fallback (void)
|
|
|
1587 {
|
|
|
1588 struct gomp_task_icv *icv = gomp_icv (false);
|
|
|
1589 if (icv->target_data)
|
|
|
1590 {
|
|
|
1591 /* Even when doing a host fallback, if there are any active
|
|
|
1592 #pragma omp target data constructs, need to remember the
|
|
|
1593 new #pragma omp target data, otherwise GOMP_target_end_data
|
|
|
1594 would get out of sync. */
|
|
|
1595 struct target_mem_desc *tgt
|
|
|
1596 = gomp_map_vars (NULL, 0, NULL, NULL, NULL, NULL, false,
|
|
|
1597 GOMP_MAP_VARS_DATA);
|
|
|
1598 tgt->prev = icv->target_data;
|
|
|
1599 icv->target_data = tgt;
|
|
|
1600 }
|
|
|
1601 }
|
|
|
1602
|
|
|
1603 void
|
|
|
1604 GOMP_target_data (int device, const void *unused, size_t mapnum,
|
|
|
1605 void **hostaddrs, size_t *sizes, unsigned char *kinds)
|
|
|
1606 {
|
|
|
1607 struct gomp_device_descr *devicep = resolve_device (device);
|
|
|
1608
|
|
|
1609 if (devicep == NULL
|
|
|
1610 || !(devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400)
|
|
|
1611 || (devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM))
|
|
|
1612 return gomp_target_data_fallback ();
|
|
|
1613
|
|
|
1614 struct target_mem_desc *tgt
|
|
|
1615 = gomp_map_vars (devicep, mapnum, hostaddrs, NULL, sizes, kinds, false,
|
|
|
1616 GOMP_MAP_VARS_DATA);
|
|
|
1617 struct gomp_task_icv *icv = gomp_icv (true);
|
|
|
1618 tgt->prev = icv->target_data;
|
|
|
1619 icv->target_data = tgt;
|
|
|
1620 }
|
|
|
1621
|
|
|
1622 void
|
|
|
1623 GOMP_target_data_ext (int device, size_t mapnum, void **hostaddrs,
|
|
|
1624 size_t *sizes, unsigned short *kinds)
|
|
|
1625 {
|
|
|
1626 struct gomp_device_descr *devicep = resolve_device (device);
|
|
|
1627
|
|
|
1628 if (devicep == NULL
|
|
|
1629 || !(devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400)
|
|
|
1630 || devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
|
|
|
1631 return gomp_target_data_fallback ();
|
|
|
1632
|
|
|
1633 struct target_mem_desc *tgt
|
|
|
1634 = gomp_map_vars (devicep, mapnum, hostaddrs, NULL, sizes, kinds, true,
|
|
|
1635 GOMP_MAP_VARS_DATA);
|
|
|
1636 struct gomp_task_icv *icv = gomp_icv (true);
|
|
|
1637 tgt->prev = icv->target_data;
|
|
|
1638 icv->target_data = tgt;
|
|
|
1639 }
|
|
|
1640
|
|
|
1641 void
|
|
|
1642 GOMP_target_end_data (void)
|
|
|
1643 {
|
|
|
1644 struct gomp_task_icv *icv = gomp_icv (false);
|
|
|
1645 if (icv->target_data)
|
|
|
1646 {
|
|
|
1647 struct target_mem_desc *tgt = icv->target_data;
|
|
|
1648 icv->target_data = tgt->prev;
|
|
|
1649 gomp_unmap_vars (tgt, true);
|
|
|
1650 }
|
|
|
1651 }
|
|
|
1652
|
|
|
1653 void
|
|
|
1654 GOMP_target_update (int device, const void *unused, size_t mapnum,
|
|
|
1655 void **hostaddrs, size_t *sizes, unsigned char *kinds)
|
|
|
1656 {
|
|
|
1657 struct gomp_device_descr *devicep = resolve_device (device);
|
|
|
1658
|
|
|
1659 if (devicep == NULL
|
|
|
1660 || !(devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400)
|
|
|
1661 || devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
|
|
|
1662 return;
|
|
|
1663
|
|
|
1664 gomp_update (devicep, mapnum, hostaddrs, sizes, kinds, false);
|
|
|
1665 }
|
|
|
1666
|
|
|
1667 void
|
|
|
1668 GOMP_target_update_ext (int device, size_t mapnum, void **hostaddrs,
|
|
|
1669 size_t *sizes, unsigned short *kinds,
|
|
|
1670 unsigned int flags, void **depend)
|
|
|
1671 {
|
|
|
1672 struct gomp_device_descr *devicep = resolve_device (device);
|
|
|
1673
|
|
|
1674 /* If there are depend clauses, but nowait is not present,
|
|
|
1675 block the parent task until the dependencies are resolved
|
|
|
1676 and then just continue with the rest of the function as if it
|
|
|
1677 is a merged task. Until we are able to schedule task during
|
|
|
1678 variable mapping or unmapping, ignore nowait if depend clauses
|
|
|
1679 are not present. */
|
|
|
1680 if (depend != NULL)
|
|
|
1681 {
|
|
|
1682 struct gomp_thread *thr = gomp_thread ();
|
|
|
1683 if (thr->task && thr->task->depend_hash)
|
|
|
1684 {
|
|
|
1685 if ((flags & GOMP_TARGET_FLAG_NOWAIT)
|
|
|
1686 && thr->ts.team
|
|
|
1687 && !thr->task->final_task)
|
|
|
1688 {
|
|
|
1689 if (gomp_create_target_task (devicep, (void (*) (void *)) NULL,
|
|
|
1690 mapnum, hostaddrs, sizes, kinds,
|
|
|
1691 flags | GOMP_TARGET_FLAG_UPDATE,
|
|
|
1692 depend, NULL, GOMP_TARGET_TASK_DATA))
|
|
|
1693 return;
|
|
|
1694 }
|
|
|
1695 else
|
|
|
1696 {
|
|
|
1697 struct gomp_team *team = thr->ts.team;
|
|
|
1698 /* If parallel or taskgroup has been cancelled, don't start new
|
|
|
1699 tasks. */
|
|
|
1700 if (team
|
|
|
1701 && (gomp_team_barrier_cancelled (&team->barrier)
|
|
|
1702 || (thr->task->taskgroup
|
|
|
1703 && thr->task->taskgroup->cancelled)))
|
|
|
1704 return;
|
|
|
1705
|
|
|
1706 gomp_task_maybe_wait_for_dependencies (depend);
|
|
|
1707 }
|
|
|
1708 }
|
|
|
1709 }
|
|
|
1710
|
|
|
1711 if (devicep == NULL
|
|
|
1712 || !(devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400)
|
|
|
1713 || devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
|
|
|
1714 return;
|
|
|
1715
|
|
|
1716 struct gomp_thread *thr = gomp_thread ();
|
|
|
1717 struct gomp_team *team = thr->ts.team;
|
|
|
1718 /* If parallel or taskgroup has been cancelled, don't start new tasks. */
|
|
|
1719 if (team
|
|
|
1720 && (gomp_team_barrier_cancelled (&team->barrier)
|
|
|
1721 || (thr->task->taskgroup && thr->task->taskgroup->cancelled)))
|
|
|
1722 return;
|
|
|
1723
|
|
|
1724 gomp_update (devicep, mapnum, hostaddrs, sizes, kinds, true);
|
|
|
1725 }
|
|
|
1726
|
|
|
1727 static void
|
|
|
1728 gomp_exit_data (struct gomp_device_descr *devicep, size_t mapnum,
|
|
|
1729 void **hostaddrs, size_t *sizes, unsigned short *kinds)
|
|
|
1730 {
|
|
|
1731 const int typemask = 0xff;
|
|
|
1732 size_t i;
|
|
|
1733 gomp_mutex_lock (&devicep->lock);
|
|
|
1734 if (devicep->state == GOMP_DEVICE_FINALIZED)
|
|
|
1735 {
|
|
|
1736 gomp_mutex_unlock (&devicep->lock);
|
|
|
1737 return;
|
|
|
1738 }
|
|
|
1739
|
|
|
1740 for (i = 0; i < mapnum; i++)
|
|
|
1741 {
|
|
|
1742 struct splay_tree_key_s cur_node;
|
|
|
1743 unsigned char kind = kinds[i] & typemask;
|
|
|
1744 switch (kind)
|
|
|
1745 {
|
|
|
1746 case GOMP_MAP_FROM:
|
|
|
1747 case GOMP_MAP_ALWAYS_FROM:
|
|
|
1748 case GOMP_MAP_DELETE:
|
|
|
1749 case GOMP_MAP_RELEASE:
|
|
|
1750 case GOMP_MAP_ZERO_LEN_ARRAY_SECTION:
|
|
|
1751 case GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION:
|
|
|
1752 cur_node.host_start = (uintptr_t) hostaddrs[i];
|
|
|
1753 cur_node.host_end = cur_node.host_start + sizes[i];
|
|
|
1754 splay_tree_key k = (kind == GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION
|
|
|
1755 || kind == GOMP_MAP_ZERO_LEN_ARRAY_SECTION)
|
|
|
1756 ? gomp_map_0len_lookup (&devicep->mem_map, &cur_node)
|
|
|
1757 : splay_tree_lookup (&devicep->mem_map, &cur_node);
|
|
|
1758 if (!k)
|
|
|
1759 continue;
|
|
|
1760
|
|
|
1761 if (k->refcount > 0 && k->refcount != REFCOUNT_INFINITY)
|
|
|
1762 k->refcount--;
|
|
|
1763 if ((kind == GOMP_MAP_DELETE
|
|
|
1764 || kind == GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION)
|
|
|
1765 && k->refcount != REFCOUNT_INFINITY)
|
|
|
1766 k->refcount = 0;
|
|
|
1767
|
|
|
1768 if ((kind == GOMP_MAP_FROM && k->refcount == 0)
|
|
|
1769 || kind == GOMP_MAP_ALWAYS_FROM)
|
|
|
1770 gomp_copy_dev2host (devicep, (void *) cur_node.host_start,
|
|
|
1771 (void *) (k->tgt->tgt_start + k->tgt_offset
|
|
|
1772 + cur_node.host_start
|
|
|
1773 - k->host_start),
|
|
|
1774 cur_node.host_end - cur_node.host_start);
|
|
|
1775 if (k->refcount == 0)
|
|
|
1776 {
|
|
|
1777 splay_tree_remove (&devicep->mem_map, k);
|
|
|
1778 if (k->link_key)
|
|
|
1779 splay_tree_insert (&devicep->mem_map,
|
|
|
1780 (splay_tree_node) k->link_key);
|
|
|
1781 if (k->tgt->refcount > 1)
|
|
|
1782 k->tgt->refcount--;
|
|
|
1783 else
|
|
|
1784 gomp_unmap_tgt (k->tgt);
|
|
|
1785 }
|
|
|
1786
|
|
|
1787 break;
|
|
|
1788 default:
|
|
|
1789 gomp_mutex_unlock (&devicep->lock);
|
|
|
1790 gomp_fatal ("GOMP_target_enter_exit_data unhandled kind 0x%.2x",
|
|
|
1791 kind);
|
|
|
1792 }
|
|
|
1793 }
|
|
|
1794
|
|
|
1795 gomp_mutex_unlock (&devicep->lock);
|
|
|
1796 }
|
|
|
1797
|
|
|
1798 void
|
|
|
1799 GOMP_target_enter_exit_data (int device, size_t mapnum, void **hostaddrs,
|
|
|
1800 size_t *sizes, unsigned short *kinds,
|
|
|
1801 unsigned int flags, void **depend)
|
|
|
1802 {
|
|
|
1803 struct gomp_device_descr *devicep = resolve_device (device);
|
|
|
1804
|
|
|
1805 /* If there are depend clauses, but nowait is not present,
|
|
|
1806 block the parent task until the dependencies are resolved
|
|
|
1807 and then just continue with the rest of the function as if it
|
|
|
1808 is a merged task. Until we are able to schedule task during
|
|
|
1809 variable mapping or unmapping, ignore nowait if depend clauses
|
|
|
1810 are not present. */
|
|
|
1811 if (depend != NULL)
|
|
|
1812 {
|
|
|
1813 struct gomp_thread *thr = gomp_thread ();
|
|
|
1814 if (thr->task && thr->task->depend_hash)
|
|
|
1815 {
|
|
|
1816 if ((flags & GOMP_TARGET_FLAG_NOWAIT)
|
|
|
1817 && thr->ts.team
|
|
|
1818 && !thr->task->final_task)
|
|
|
1819 {
|
|
|
1820 if (gomp_create_target_task (devicep, (void (*) (void *)) NULL,
|
|
|
1821 mapnum, hostaddrs, sizes, kinds,
|
|
|
1822 flags, depend, NULL,
|
|
|
1823 GOMP_TARGET_TASK_DATA))
|
|
|
1824 return;
|
|
|
1825 }
|
|
|
1826 else
|
|
|
1827 {
|
|
|
1828 struct gomp_team *team = thr->ts.team;
|
|
|
1829 /* If parallel or taskgroup has been cancelled, don't start new
|
|
|
1830 tasks. */
|
|
|
1831 if (team
|
|
|
1832 && (gomp_team_barrier_cancelled (&team->barrier)
|
|
|
1833 || (thr->task->taskgroup
|
|
|
1834 && thr->task->taskgroup->cancelled)))
|
|
|
1835 return;
|
|
|
1836
|
|
|
1837 gomp_task_maybe_wait_for_dependencies (depend);
|
|
|
1838 }
|
|
|
1839 }
|
|
|
1840 }
|
|
|
1841
|
|
|
1842 if (devicep == NULL
|
|
|
1843 || !(devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400)
|
|
|
1844 || devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
|
|
|
1845 return;
|
|
|
1846
|
|
|
1847 struct gomp_thread *thr = gomp_thread ();
|
|
|
1848 struct gomp_team *team = thr->ts.team;
|
|
|
1849 /* If parallel or taskgroup has been cancelled, don't start new tasks. */
|
|
|
1850 if (team
|
|
|
1851 && (gomp_team_barrier_cancelled (&team->barrier)
|
|
|
1852 || (thr->task->taskgroup && thr->task->taskgroup->cancelled)))
|
|
|
1853 return;
|
|
|
1854
|
|
|
1855 size_t i;
|
|
|
1856 if ((flags & GOMP_TARGET_FLAG_EXIT_DATA) == 0)
|
|
|
1857 for (i = 0; i < mapnum; i++)
|
|
|
1858 if ((kinds[i] & 0xff) == GOMP_MAP_STRUCT)
|
|
|
1859 {
|
|
|
1860 gomp_map_vars (devicep, sizes[i] + 1, &hostaddrs[i], NULL, &sizes[i],
|
|
|
1861 &kinds[i], true, GOMP_MAP_VARS_ENTER_DATA);
|
|
|
1862 i += sizes[i];
|
|
|
1863 }
|
|
|
1864 else
|
|
|
1865 gomp_map_vars (devicep, 1, &hostaddrs[i], NULL, &sizes[i], &kinds[i],
|
|
|
1866 true, GOMP_MAP_VARS_ENTER_DATA);
|
|
|
1867 else
|
|
|
1868 gomp_exit_data (devicep, mapnum, hostaddrs, sizes, kinds);
|
|
|
1869 }
|
|
|
1870
|
|
|
1871 bool
|
|
|
1872 gomp_target_task_fn (void *data)
|
|
|
1873 {
|
|
|
1874 struct gomp_target_task *ttask = (struct gomp_target_task *) data;
|
|
|
1875 struct gomp_device_descr *devicep = ttask->devicep;
|
|
|
1876
|
|
|
1877 if (ttask->fn != NULL)
|
|
|
1878 {
|
|
|
1879 void *fn_addr;
|
|
|
1880 if (devicep == NULL
|
|
|
1881 || !(devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400)
|
|
|
1882 || !(fn_addr = gomp_get_target_fn_addr (devicep, ttask->fn))
|
|
|
1883 || (devicep->can_run_func && !devicep->can_run_func (fn_addr)))
|
|
|
1884 {
|
|
|
1885 ttask->state = GOMP_TARGET_TASK_FALLBACK;
|
|
|
1886 gomp_target_fallback (ttask->fn, ttask->hostaddrs);
|
|
|
1887 return false;
|
|
|
1888 }
|
|
|
1889
|
|
|
1890 if (ttask->state == GOMP_TARGET_TASK_FINISHED)
|
|
|
1891 {
|
|
|
1892 if (ttask->tgt)
|
|
|
1893 gomp_unmap_vars (ttask->tgt, true);
|
|
|
1894 return false;
|
|
|
1895 }
|
|
|
1896
|
|
|
1897 void *actual_arguments;
|
|
|
1898 if (devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
|
|
|
1899 {
|
|
|
1900 ttask->tgt = NULL;
|
|
|
1901 actual_arguments = ttask->hostaddrs;
|
|
|
1902 }
|
|
|
1903 else
|
|
|
1904 {
|
|
|
1905 ttask->tgt = gomp_map_vars (devicep, ttask->mapnum, ttask->hostaddrs,
|
|
|
1906 NULL, ttask->sizes, ttask->kinds, true,
|
|
|
1907 GOMP_MAP_VARS_TARGET);
|
|
|
1908 actual_arguments = (void *) ttask->tgt->tgt_start;
|
|
|
1909 }
|
|
|
1910 ttask->state = GOMP_TARGET_TASK_READY_TO_RUN;
|
|
|
1911
|
|
|
1912 devicep->async_run_func (devicep->target_id, fn_addr, actual_arguments,
|
|
|
1913 ttask->args, (void *) ttask);
|
|
|
1914 return true;
|
|
|
1915 }
|
|
|
1916 else if (devicep == NULL
|
|
|
1917 || !(devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400)
|
|
|
1918 || devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
|
|
|
1919 return false;
|
|
|
1920
|
|
|
1921 size_t i;
|
|
|
1922 if (ttask->flags & GOMP_TARGET_FLAG_UPDATE)
|
|
|
1923 gomp_update (devicep, ttask->mapnum, ttask->hostaddrs, ttask->sizes,
|
|
|
1924 ttask->kinds, true);
|
|
|
1925 else if ((ttask->flags & GOMP_TARGET_FLAG_EXIT_DATA) == 0)
|
|
|
1926 for (i = 0; i < ttask->mapnum; i++)
|
|
|
1927 if ((ttask->kinds[i] & 0xff) == GOMP_MAP_STRUCT)
|
|
|
1928 {
|
|
|
1929 gomp_map_vars (devicep, ttask->sizes[i] + 1, &ttask->hostaddrs[i],
|
|
|
1930 NULL, &ttask->sizes[i], &ttask->kinds[i], true,
|
|
|
1931 GOMP_MAP_VARS_ENTER_DATA);
|
|
|
1932 i += ttask->sizes[i];
|
|
|
1933 }
|
|
|
1934 else
|
|
|
1935 gomp_map_vars (devicep, 1, &ttask->hostaddrs[i], NULL, &ttask->sizes[i],
|
|
|
1936 &ttask->kinds[i], true, GOMP_MAP_VARS_ENTER_DATA);
|
|
|
1937 else
|
|
|
1938 gomp_exit_data (devicep, ttask->mapnum, ttask->hostaddrs, ttask->sizes,
|
|
|
1939 ttask->kinds);
|
|
|
1940 return false;
|
|
|
1941 }
|
|
|
1942
|
|
|
1943 void
|
|
|
1944 GOMP_teams (unsigned int num_teams, unsigned int thread_limit)
|
|
|
1945 {
|
|
|
1946 if (thread_limit)
|
|
|
1947 {
|
|
|
1948 struct gomp_task_icv *icv = gomp_icv (true);
|
|
|
1949 icv->thread_limit_var
|
|
|
1950 = thread_limit > INT_MAX ? UINT_MAX : thread_limit;
|
|
|
1951 }
|
|
|
1952 (void) num_teams;
|
|
|
1953 }
|
|
|
1954
|
|
|
1955 void *
|
|
|
1956 omp_target_alloc (size_t size, int device_num)
|
|
|
1957 {
|
|
|
1958 if (device_num == GOMP_DEVICE_HOST_FALLBACK)
|
|
|
1959 return malloc (size);
|
|
|
1960
|
|
|
1961 if (device_num < 0)
|
|
|
1962 return NULL;
|
|
|
1963
|
|
|
1964 struct gomp_device_descr *devicep = resolve_device (device_num);
|
|
|
1965 if (devicep == NULL)
|
|
|
1966 return NULL;
|
|
|
1967
|
|
|
1968 if (!(devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400)
|
|
|
1969 || devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
|
|
|
1970 return malloc (size);
|
|
|
1971
|
|
|
1972 gomp_mutex_lock (&devicep->lock);
|
|
|
1973 void *ret = devicep->alloc_func (devicep->target_id, size);
|
|
|
1974 gomp_mutex_unlock (&devicep->lock);
|
|
|
1975 return ret;
|
|
|
1976 }
|
|
|
1977
|
|
|
1978 void
|
|
|
1979 omp_target_free (void *device_ptr, int device_num)
|
|
|
1980 {
|
|
|
1981 if (device_ptr == NULL)
|
|
|
1982 return;
|
|
|
1983
|
|
|
1984 if (device_num == GOMP_DEVICE_HOST_FALLBACK)
|
|
|
1985 {
|
|
|
1986 free (device_ptr);
|
|
|
1987 return;
|
|
|
1988 }
|
|
|
1989
|
|
|
1990 if (device_num < 0)
|
|
|
1991 return;
|
|
|
1992
|
|
|
1993 struct gomp_device_descr *devicep = resolve_device (device_num);
|
|
|
1994 if (devicep == NULL)
|
|
|
1995 return;
|
|
|
1996
|
|
|
1997 if (!(devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400)
|
|
|
1998 || devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
|
|
|
1999 {
|
|
|
2000 free (device_ptr);
|
|
|
2001 return;
|
|
|
2002 }
|
|
|
2003
|
|
|
2004 gomp_mutex_lock (&devicep->lock);
|
|
|
2005 gomp_free_device_memory (devicep, device_ptr);
|
|
|
2006 gomp_mutex_unlock (&devicep->lock);
|
|
|
2007 }
|
|
|
2008
|
|
|
2009 int
|
|
|
2010 omp_target_is_present (void *ptr, int device_num)
|
|
|
2011 {
|
|
|
2012 if (ptr == NULL)
|
|
|
2013 return 1;
|
|
|
2014
|
|
|
2015 if (device_num == GOMP_DEVICE_HOST_FALLBACK)
|
|
|
2016 return 1;
|
|
|
2017
|
|
|
2018 if (device_num < 0)
|
|
|
2019 return 0;
|
|
|
2020
|
|
|
2021 struct gomp_device_descr *devicep = resolve_device (device_num);
|
|
|
2022 if (devicep == NULL)
|
|
|
2023 return 0;
|
|
|
2024
|
|
|
2025 if (!(devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400)
|
|
|
2026 || devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
|
|
|
2027 return 1;
|
|
|
2028
|
|
|
2029 gomp_mutex_lock (&devicep->lock);
|
|
|
2030 struct splay_tree_s *mem_map = &devicep->mem_map;
|
|
|
2031 struct splay_tree_key_s cur_node;
|
|
|
2032
|
|
|
2033 cur_node.host_start = (uintptr_t) ptr;
|
|
|
2034 cur_node.host_end = cur_node.host_start;
|
|
|
2035 splay_tree_key n = gomp_map_0len_lookup (mem_map, &cur_node);
|
|
|
2036 int ret = n != NULL;
|
|
|
2037 gomp_mutex_unlock (&devicep->lock);
|
|
|
2038 return ret;
|
|
|
2039 }
|
|
|
2040
|
|
|
2041 int
|
|
|
2042 omp_target_memcpy (void *dst, void *src, size_t length, size_t dst_offset,
|
|
|
2043 size_t src_offset, int dst_device_num, int src_device_num)
|
|
|
2044 {
|
|
|
2045 struct gomp_device_descr *dst_devicep = NULL, *src_devicep = NULL;
|
|
|
2046 bool ret;
|
|
|
2047
|
|
|
2048 if (dst_device_num != GOMP_DEVICE_HOST_FALLBACK)
|
|
|
2049 {
|
|
|
2050 if (dst_device_num < 0)
|
|
|
2051 return EINVAL;
|
|
|
2052
|
|
|
2053 dst_devicep = resolve_device (dst_device_num);
|
|
|
2054 if (dst_devicep == NULL)
|
|
|
2055 return EINVAL;
|
|
|
2056
|
|
|
2057 if (!(dst_devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400)
|
|
|
2058 || dst_devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
|
|
|
2059 dst_devicep = NULL;
|
|
|
2060 }
|
|
|
2061 if (src_device_num != GOMP_DEVICE_HOST_FALLBACK)
|
|
|
2062 {
|
|
|
2063 if (src_device_num < 0)
|
|
|
2064 return EINVAL;
|
|
|
2065
|
|
|
2066 src_devicep = resolve_device (src_device_num);
|
|
|
2067 if (src_devicep == NULL)
|
|
|
2068 return EINVAL;
|
|
|
2069
|
|
|
2070 if (!(src_devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400)
|
|
|
2071 || src_devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
|
|
|
2072 src_devicep = NULL;
|
|
|
2073 }
|
|
|
2074 if (src_devicep == NULL && dst_devicep == NULL)
|
|
|
2075 {
|
|
|
2076 memcpy ((char *) dst + dst_offset, (char *) src + src_offset, length);
|
|
|
2077 return 0;
|
|
|
2078 }
|
|
|
2079 if (src_devicep == NULL)
|
|
|
2080 {
|
|
|
2081 gomp_mutex_lock (&dst_devicep->lock);
|
|
|
2082 ret = dst_devicep->host2dev_func (dst_devicep->target_id,
|
|
|
2083 (char *) dst + dst_offset,
|
|
|
2084 (char *) src + src_offset, length);
|
|
|
2085 gomp_mutex_unlock (&dst_devicep->lock);
|
|
|
2086 return (ret ? 0 : EINVAL);
|
|
|
2087 }
|
|
|
2088 if (dst_devicep == NULL)
|
|
|
2089 {
|
|
|
2090 gomp_mutex_lock (&src_devicep->lock);
|
|
|
2091 ret = src_devicep->dev2host_func (src_devicep->target_id,
|
|
|
2092 (char *) dst + dst_offset,
|
|
|
2093 (char *) src + src_offset, length);
|
|
|
2094 gomp_mutex_unlock (&src_devicep->lock);
|
|
|
2095 return (ret ? 0 : EINVAL);
|
|
|
2096 }
|
|
|
2097 if (src_devicep == dst_devicep)
|
|
|
2098 {
|
|
|
2099 gomp_mutex_lock (&src_devicep->lock);
|
|
|
2100 ret = src_devicep->dev2dev_func (src_devicep->target_id,
|
|
|
2101 (char *) dst + dst_offset,
|
|
|
2102 (char *) src + src_offset, length);
|
|
|
2103 gomp_mutex_unlock (&src_devicep->lock);
|
|
|
2104 return (ret ? 0 : EINVAL);
|
|
|
2105 }
|
|
|
2106 return EINVAL;
|
|
|
2107 }
|
|
|
2108
|
|
|
2109 static int
|
|
|
2110 omp_target_memcpy_rect_worker (void *dst, void *src, size_t element_size,
|
|
|
2111 int num_dims, const size_t *volume,
|
|
|
2112 const size_t *dst_offsets,
|
|
|
2113 const size_t *src_offsets,
|
|
|
2114 const size_t *dst_dimensions,
|
|
|
2115 const size_t *src_dimensions,
|
|
|
2116 struct gomp_device_descr *dst_devicep,
|
|
|
2117 struct gomp_device_descr *src_devicep)
|
|
|
2118 {
|
|
|
2119 size_t dst_slice = element_size;
|
|
|
2120 size_t src_slice = element_size;
|
|
|
2121 size_t j, dst_off, src_off, length;
|
|
|
2122 int i, ret;
|
|
|
2123
|
|
|
2124 if (num_dims == 1)
|
|
|
2125 {
|
|
|
2126 if (__builtin_mul_overflow (element_size, volume[0], &length)
|
|
|
2127 || __builtin_mul_overflow (element_size, dst_offsets[0], &dst_off)
|
|
|
2128 || __builtin_mul_overflow (element_size, src_offsets[0], &src_off))
|
|
|
2129 return EINVAL;
|
|
|
2130 if (dst_devicep == NULL && src_devicep == NULL)
|
|
|
2131 {
|
|
|
2132 memcpy ((char *) dst + dst_off, (char *) src + src_off, length);
|
|
|
2133 ret = 1;
|
|
|
2134 }
|
|
|
2135 else if (src_devicep == NULL)
|
|
|
2136 ret = dst_devicep->host2dev_func (dst_devicep->target_id,
|
|
|
2137 (char *) dst + dst_off,
|
|
|
2138 (char *) src + src_off, length);
|
|
|
2139 else if (dst_devicep == NULL)
|
|
|
2140 ret = src_devicep->dev2host_func (src_devicep->target_id,
|
|
|
2141 (char *) dst + dst_off,
|
|
|
2142 (char *) src + src_off, length);
|
|
|
2143 else if (src_devicep == dst_devicep)
|
|
|
2144 ret = src_devicep->dev2dev_func (src_devicep->target_id,
|
|
|
2145 (char *) dst + dst_off,
|
|
|
2146 (char *) src + src_off, length);
|
|
|
2147 else
|
|
|
2148 ret = 0;
|
|
|
2149 return ret ? 0 : EINVAL;
|
|
|
2150 }
|
|
|
2151
|
|
|
2152 /* FIXME: it would be nice to have some plugin function to handle
|
|
|
2153 num_dims == 2 and num_dims == 3 more efficiently. Larger ones can
|
|
|
2154 be handled in the generic recursion below, and for host-host it
|
|
|
2155 should be used even for any num_dims >= 2. */
|
|
|
2156
|
|
|
2157 for (i = 1; i < num_dims; i++)
|
|
|
2158 if (__builtin_mul_overflow (dst_slice, dst_dimensions[i], &dst_slice)
|
|
|
2159 || __builtin_mul_overflow (src_slice, src_dimensions[i], &src_slice))
|
|
|
2160 return EINVAL;
|
|
|
2161 if (__builtin_mul_overflow (dst_slice, dst_offsets[0], &dst_off)
|
|
|
2162 || __builtin_mul_overflow (src_slice, src_offsets[0], &src_off))
|
|
|
2163 return EINVAL;
|
|
|
2164 for (j = 0; j < volume[0]; j++)
|
|
|
2165 {
|
|
|
2166 ret = omp_target_memcpy_rect_worker ((char *) dst + dst_off,
|
|
|
2167 (char *) src + src_off,
|
|
|
2168 element_size, num_dims - 1,
|
|
|
2169 volume + 1, dst_offsets + 1,
|
|
|
2170 src_offsets + 1, dst_dimensions + 1,
|
|
|
2171 src_dimensions + 1, dst_devicep,
|
|
|
2172 src_devicep);
|
|
|
2173 if (ret)
|
|
|
2174 return ret;
|
|
|
2175 dst_off += dst_slice;
|
|
|
2176 src_off += src_slice;
|
|
|
2177 }
|
|
|
2178 return 0;
|
|
|
2179 }
|
|
|
2180
|
|
|
2181 int
|
|
|
2182 omp_target_memcpy_rect (void *dst, void *src, size_t element_size,
|
|
|
2183 int num_dims, const size_t *volume,
|
|
|
2184 const size_t *dst_offsets,
|
|
|
2185 const size_t *src_offsets,
|
|
|
2186 const size_t *dst_dimensions,
|
|
|
2187 const size_t *src_dimensions,
|
|
|
2188 int dst_device_num, int src_device_num)
|
|
|
2189 {
|
|
|
2190 struct gomp_device_descr *dst_devicep = NULL, *src_devicep = NULL;
|
|
|
2191
|
|
|
2192 if (!dst && !src)
|
|
|
2193 return INT_MAX;
|
|
|
2194
|
|
|
2195 if (dst_device_num != GOMP_DEVICE_HOST_FALLBACK)
|
|
|
2196 {
|
|
|
2197 if (dst_device_num < 0)
|
|
|
2198 return EINVAL;
|
|
|
2199
|
|
|
2200 dst_devicep = resolve_device (dst_device_num);
|
|
|
2201 if (dst_devicep == NULL)
|
|
|
2202 return EINVAL;
|
|
|
2203
|
|
|
2204 if (!(dst_devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400)
|
|
|
2205 || dst_devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
|
|
|
2206 dst_devicep = NULL;
|
|
|
2207 }
|
|
|
2208 if (src_device_num != GOMP_DEVICE_HOST_FALLBACK)
|
|
|
2209 {
|
|
|
2210 if (src_device_num < 0)
|
|
|
2211 return EINVAL;
|
|
|
2212
|
|
|
2213 src_devicep = resolve_device (src_device_num);
|
|
|
2214 if (src_devicep == NULL)
|
|
|
2215 return EINVAL;
|
|
|
2216
|
|
|
2217 if (!(src_devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400)
|
|
|
2218 || src_devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
|
|
|
2219 src_devicep = NULL;
|
|
|
2220 }
|
|
|
2221
|
|
|
2222 if (src_devicep != NULL && dst_devicep != NULL && src_devicep != dst_devicep)
|
|
|
2223 return EINVAL;
|
|
|
2224
|
|
|
2225 if (src_devicep)
|
|
|
2226 gomp_mutex_lock (&src_devicep->lock);
|
|
|
2227 else if (dst_devicep)
|
|
|
2228 gomp_mutex_lock (&dst_devicep->lock);
|
|
|
2229 int ret = omp_target_memcpy_rect_worker (dst, src, element_size, num_dims,
|
|
|
2230 volume, dst_offsets, src_offsets,
|
|
|
2231 dst_dimensions, src_dimensions,
|
|
|
2232 dst_devicep, src_devicep);
|
|
|
2233 if (src_devicep)
|
|
|
2234 gomp_mutex_unlock (&src_devicep->lock);
|
|
|
2235 else if (dst_devicep)
|
|
|
2236 gomp_mutex_unlock (&dst_devicep->lock);
|
|
|
2237 return ret;
|
|
|
2238 }
|
|
|
2239
|
|
|
2240 int
|
|
|
2241 omp_target_associate_ptr (void *host_ptr, void *device_ptr, size_t size,
|
|
|
2242 size_t device_offset, int device_num)
|
|
|
2243 {
|
|
|
2244 if (device_num == GOMP_DEVICE_HOST_FALLBACK)
|
|
|
2245 return EINVAL;
|
|
|
2246
|
|
|
2247 if (device_num < 0)
|
|
|
2248 return EINVAL;
|
|
|
2249
|
|
|
2250 struct gomp_device_descr *devicep = resolve_device (device_num);
|
|
|
2251 if (devicep == NULL)
|
|
|
2252 return EINVAL;
|
|
|
2253
|
|
|
2254 if (!(devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400)
|
|
|
2255 || devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
|
|
|
2256 return EINVAL;
|
|
|
2257
|
|
|
2258 gomp_mutex_lock (&devicep->lock);
|
|
|
2259
|
|
|
2260 struct splay_tree_s *mem_map = &devicep->mem_map;
|
|
|
2261 struct splay_tree_key_s cur_node;
|
|
|
2262 int ret = EINVAL;
|
|
|
2263
|
|
|
2264 cur_node.host_start = (uintptr_t) host_ptr;
|
|
|
2265 cur_node.host_end = cur_node.host_start + size;
|
|
|
2266 splay_tree_key n = gomp_map_lookup (mem_map, &cur_node);
|
|
|
2267 if (n)
|
|
|
2268 {
|
|
|
2269 if (n->tgt->tgt_start + n->tgt_offset
|
|
|
2270 == (uintptr_t) device_ptr + device_offset
|
|
|
2271 && n->host_start <= cur_node.host_start
|
|
|
2272 && n->host_end >= cur_node.host_end)
|
|
|
2273 ret = 0;
|
|
|
2274 }
|
|
|
2275 else
|
|
|
2276 {
|
|
|
2277 struct target_mem_desc *tgt = gomp_malloc (sizeof (*tgt));
|
|
|
2278 tgt->array = gomp_malloc (sizeof (*tgt->array));
|
|
|
2279 tgt->refcount = 1;
|
|
|
2280 tgt->tgt_start = 0;
|
|
|
2281 tgt->tgt_end = 0;
|
|
|
2282 tgt->to_free = NULL;
|
|
|
2283 tgt->prev = NULL;
|
|
|
2284 tgt->list_count = 0;
|
|
|
2285 tgt->device_descr = devicep;
|
|
|
2286 splay_tree_node array = tgt->array;
|
|
|
2287 splay_tree_key k = &array->key;
|
|
|
2288 k->host_start = cur_node.host_start;
|
|
|
2289 k->host_end = cur_node.host_end;
|
|
|
2290 k->tgt = tgt;
|
|
|
2291 k->tgt_offset = (uintptr_t) device_ptr + device_offset;
|
|
|
2292 k->refcount = REFCOUNT_INFINITY;
|
|
|
2293 array->left = NULL;
|
|
|
2294 array->right = NULL;
|
|
|
2295 splay_tree_insert (&devicep->mem_map, array);
|
|
|
2296 ret = 0;
|
|
|
2297 }
|
|
|
2298 gomp_mutex_unlock (&devicep->lock);
|
|
|
2299 return ret;
|
|
|
2300 }
|
|
|
2301
|
|
|
2302 int
|
|
|
2303 omp_target_disassociate_ptr (void *ptr, int device_num)
|
|
|
2304 {
|
|
|
2305 if (device_num == GOMP_DEVICE_HOST_FALLBACK)
|
|
|
2306 return EINVAL;
|
|
|
2307
|
|
|
2308 if (device_num < 0)
|
|
|
2309 return EINVAL;
|
|
|
2310
|
|
|
2311 struct gomp_device_descr *devicep = resolve_device (device_num);
|
|
|
2312 if (devicep == NULL)
|
|
|
2313 return EINVAL;
|
|
|
2314
|
|
|
2315 if (!(devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400))
|
|
|
2316 return EINVAL;
|
|
|
2317
|
|
|
2318 gomp_mutex_lock (&devicep->lock);
|
|
|
2319
|
|
|
2320 struct splay_tree_s *mem_map = &devicep->mem_map;
|
|
|
2321 struct splay_tree_key_s cur_node;
|
|
|
2322 int ret = EINVAL;
|
|
|
2323
|
|
|
2324 cur_node.host_start = (uintptr_t) ptr;
|
|
|
2325 cur_node.host_end = cur_node.host_start;
|
|
|
2326 splay_tree_key n = gomp_map_lookup (mem_map, &cur_node);
|
|
|
2327 if (n
|
|
|
2328 && n->host_start == cur_node.host_start
|
|
|
2329 && n->refcount == REFCOUNT_INFINITY
|
|
|
2330 && n->tgt->tgt_start == 0
|
|
|
2331 && n->tgt->to_free == NULL
|
|
|
2332 && n->tgt->refcount == 1
|
|
|
2333 && n->tgt->list_count == 0)
|
|
|
2334 {
|
|
|
2335 splay_tree_remove (&devicep->mem_map, n);
|
|
|
2336 gomp_unmap_tgt (n->tgt);
|
|
|
2337 ret = 0;
|
|
|
2338 }
|
|
|
2339
|
|
|
2340 gomp_mutex_unlock (&devicep->lock);
|
|
|
2341 return ret;
|
|
|
2342 }
|
|
|
2343
|
|
|
2344 #ifdef PLUGIN_SUPPORT
|
|
|
2345
|
|
|
2346 /* This function tries to load a plugin for DEVICE. Name of plugin is passed
|
|
|
2347 in PLUGIN_NAME.
|
|
|
2348 The handles of the found functions are stored in the corresponding fields
|
|
|
2349 of DEVICE. The function returns TRUE on success and FALSE otherwise. */
|
|
|
2350
|
|
|
2351 static bool
|
|
|
2352 gomp_load_plugin_for_device (struct gomp_device_descr *device,
|
|
|
2353 const char *plugin_name)
|
|
|
2354 {
|
|
|
2355 const char *err = NULL, *last_missing = NULL;
|
|
|
2356
|
|
|
2357 void *plugin_handle = dlopen (plugin_name, RTLD_LAZY);
|
|
|
2358 if (!plugin_handle)
|
|
|
2359 goto dl_fail;
|
|
|
2360
|
|
|
2361 /* Check if all required functions are available in the plugin and store
|
|
|
2362 their handlers. None of the symbols can legitimately be NULL,
|
|
|
2363 so we don't need to check dlerror all the time. */
|
|
|
2364 #define DLSYM(f) \
|
|
|
2365 if (!(device->f##_func = dlsym (plugin_handle, "GOMP_OFFLOAD_" #f))) \
|
|
|
2366 goto dl_fail
|
|
|
2367 /* Similar, but missing functions are not an error. Return false if
|
|
|
2368 failed, true otherwise. */
|
|
|
2369 #define DLSYM_OPT(f, n) \
|
|
|
2370 ((device->f##_func = dlsym (plugin_handle, "GOMP_OFFLOAD_" #n)) \
|
|
|
2371 || (last_missing = #n, 0))
|
|
|
2372
|
|
|
2373 DLSYM (version);
|
|
|
2374 if (device->version_func () != GOMP_VERSION)
|
|
|
2375 {
|
|
|
2376 err = "plugin version mismatch";
|
|
|
2377 goto fail;
|
|
|
2378 }
|
|
|
2379
|
|
|
2380 DLSYM (get_name);
|
|
|
2381 DLSYM (get_caps);
|
|
|
2382 DLSYM (get_type);
|
|
|
2383 DLSYM (get_num_devices);
|
|
|
2384 DLSYM (init_device);
|
|
|
2385 DLSYM (fini_device);
|
|
|
2386 DLSYM (load_image);
|
|
|
2387 DLSYM (unload_image);
|
|
|
2388 DLSYM (alloc);
|
|
|
2389 DLSYM (free);
|
|
|
2390 DLSYM (dev2host);
|
|
|
2391 DLSYM (host2dev);
|
|
|
2392 device->capabilities = device->get_caps_func ();
|
|
|
2393 if (device->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400)
|
|
|
2394 {
|
|
|
2395 DLSYM (run);
|
|
|
2396 DLSYM (async_run);
|
|
|
2397 DLSYM_OPT (can_run, can_run);
|
|
|
2398 DLSYM (dev2dev);
|
|
|
2399 }
|
|
|
2400 if (device->capabilities & GOMP_OFFLOAD_CAP_OPENACC_200)
|
|
|
2401 {
|
|
|
2402 if (!DLSYM_OPT (openacc.exec, openacc_exec)
|
|
|
2403 || !DLSYM_OPT (openacc.register_async_cleanup,
|
|
|
2404 openacc_register_async_cleanup)
|
|
|
2405 || !DLSYM_OPT (openacc.async_test, openacc_async_test)
|
|
|
2406 || !DLSYM_OPT (openacc.async_test_all, openacc_async_test_all)
|
|
|
2407 || !DLSYM_OPT (openacc.async_wait, openacc_async_wait)
|
|
|
2408 || !DLSYM_OPT (openacc.async_wait_async, openacc_async_wait_async)
|
|
|
2409 || !DLSYM_OPT (openacc.async_wait_all, openacc_async_wait_all)
|
|
|
2410 || !DLSYM_OPT (openacc.async_wait_all_async,
|
|
|
2411 openacc_async_wait_all_async)
|
|
|
2412 || !DLSYM_OPT (openacc.async_set_async, openacc_async_set_async)
|
|
|
2413 || !DLSYM_OPT (openacc.create_thread_data,
|
|
|
2414 openacc_create_thread_data)
|
|
|
2415 || !DLSYM_OPT (openacc.destroy_thread_data,
|
|
|
2416 openacc_destroy_thread_data))
|
|
|
2417 {
|
|
|
2418 /* Require all the OpenACC handlers if we have
|
|
|
2419 GOMP_OFFLOAD_CAP_OPENACC_200. */
|
|
|
2420 err = "plugin missing OpenACC handler function";
|
|
|
2421 goto fail;
|
|
|
2422 }
|
|
|
2423
|
|
|
2424 unsigned cuda = 0;
|
|
|
2425 cuda += DLSYM_OPT (openacc.cuda.get_current_device,
|
|
|
2426 openacc_cuda_get_current_device);
|
|
|
2427 cuda += DLSYM_OPT (openacc.cuda.get_current_context,
|
|
|
2428 openacc_cuda_get_current_context);
|
|
|
2429 cuda += DLSYM_OPT (openacc.cuda.get_stream, openacc_cuda_get_stream);
|
|
|
2430 cuda += DLSYM_OPT (openacc.cuda.set_stream, openacc_cuda_set_stream);
|
|
|
2431 if (cuda && cuda != 4)
|
|
|
2432 {
|
|
|
2433 /* Make sure all the CUDA functions are there if any of them are. */
|
|
|
2434 err = "plugin missing OpenACC CUDA handler function";
|
|
|
2435 goto fail;
|
|
|
2436 }
|
|
|
2437 }
|
|
|
2438 #undef DLSYM
|
|
|
2439 #undef DLSYM_OPT
|
|
|
2440
|
|
|
2441 return 1;
|
|
|
2442
|
|
|
2443 dl_fail:
|
|
|
2444 err = dlerror ();
|
|
|
2445 fail:
|
|
|
2446 gomp_error ("while loading %s: %s", plugin_name, err);
|
|
|
2447 if (last_missing)
|
|
|
2448 gomp_error ("missing function was %s", last_missing);
|
|
|
2449 if (plugin_handle)
|
|
|
2450 dlclose (plugin_handle);
|
|
|
2451
|
|
|
2452 return 0;
|
|
|
2453 }
|
|
|
2454
|
|
|
2455 /* This function finalizes all initialized devices. */
|
|
|
2456
|
|
|
2457 static void
|
|
|
2458 gomp_target_fini (void)
|
|
|
2459 {
|
|
|
2460 int i;
|
|
|
2461 for (i = 0; i < num_devices; i++)
|
|
|
2462 {
|
|
|
2463 bool ret = true;
|
|
|
2464 struct gomp_device_descr *devicep = &devices[i];
|
|
|
2465 gomp_mutex_lock (&devicep->lock);
|
|
|
2466 if (devicep->state == GOMP_DEVICE_INITIALIZED)
|
|
|
2467 {
|
|
|
2468 ret = devicep->fini_device_func (devicep->target_id);
|
|
|
2469 devicep->state = GOMP_DEVICE_FINALIZED;
|
|
|
2470 }
|
|
|
2471 gomp_mutex_unlock (&devicep->lock);
|
|
|
2472 if (!ret)
|
|
|
2473 gomp_fatal ("device finalization failed");
|
|
|
2474 }
|
|
|
2475 }
|
|
|
2476
|
|
|
2477 /* This function initializes the runtime needed for offloading.
|
|
|
2478 It parses the list of offload targets and tries to load the plugins for
|
|
|
2479 these targets. On return, the variables NUM_DEVICES and NUM_DEVICES_OPENMP
|
|
|
2480 will be set, and the array DEVICES initialized, containing descriptors for
|
|
|
2481 corresponding devices, first the GOMP_OFFLOAD_CAP_OPENMP_400 ones, follows
|
|
|
2482 by the others. */
|
|
|
2483
|
|
|
2484 static void
|
|
|
2485 gomp_target_init (void)
|
|
|
2486 {
|
|
|
2487 const char *prefix ="libgomp-plugin-";
|
|
|
2488 const char *suffix = SONAME_SUFFIX (1);
|
|
|
2489 const char *cur, *next;
|
|
|
2490 char *plugin_name;
|
|
|
2491 int i, new_num_devices;
|
|
|
2492
|
|
|
2493 num_devices = 0;
|
|
|
2494 devices = NULL;
|
|
|
2495
|
|
|
2496 cur = OFFLOAD_TARGETS;
|
|
|
2497 if (*cur)
|
|
|
2498 do
|
|
|
2499 {
|
|
|
2500 struct gomp_device_descr current_device;
|
|
|
2501
|
|
|
2502 next = strchr (cur, ',');
|
|
|
2503
|
|
|
2504 plugin_name = (char *) malloc (1 + (next ? next - cur : strlen (cur))
|
|
|
2505 + strlen (prefix) + strlen (suffix));
|
|
|
2506 if (!plugin_name)
|
|
|
2507 {
|
|
|
2508 num_devices = 0;
|
|
|
2509 break;
|
|
|
2510 }
|
|
|
2511
|
|
|
2512 strcpy (plugin_name, prefix);
|
|
|
2513 strncat (plugin_name, cur, next ? next - cur : strlen (cur));
|
|
|
2514 strcat (plugin_name, suffix);
|
|
|
2515
|
|
|
2516 if (gomp_load_plugin_for_device (¤t_device, plugin_name))
|
|
|
2517 {
|
|
|
2518 new_num_devices = current_device.get_num_devices_func ();
|
|
|
2519 if (new_num_devices >= 1)
|
|
|
2520 {
|
|
|
2521 /* Augment DEVICES and NUM_DEVICES. */
|
|
|
2522
|
|
|
2523 devices = realloc (devices, (num_devices + new_num_devices)
|
|
|
2524 * sizeof (struct gomp_device_descr));
|
|
|
2525 if (!devices)
|
|
|
2526 {
|
|
|
2527 num_devices = 0;
|
|
|
2528 free (plugin_name);
|
|
|
2529 break;
|
|
|
2530 }
|
|
|
2531
|
|
|
2532 current_device.name = current_device.get_name_func ();
|
|
|
2533 /* current_device.capabilities has already been set. */
|
|
|
2534 current_device.type = current_device.get_type_func ();
|
|
|
2535 current_device.mem_map.root = NULL;
|
|
|
2536 current_device.state = GOMP_DEVICE_UNINITIALIZED;
|
|
|
2537 current_device.openacc.data_environ = NULL;
|
|
|
2538 for (i = 0; i < new_num_devices; i++)
|
|
|
2539 {
|
|
|
2540 current_device.target_id = i;
|
|
|
2541 devices[num_devices] = current_device;
|
|
|
2542 gomp_mutex_init (&devices[num_devices].lock);
|
|
|
2543 num_devices++;
|
|
|
2544 }
|
|
|
2545 }
|
|
|
2546 }
|
|
|
2547
|
|
|
2548 free (plugin_name);
|
|
|
2549 cur = next + 1;
|
|
|
2550 }
|
|
|
2551 while (next);
|
|
|
2552
|
|
|
2553 /* In DEVICES, sort the GOMP_OFFLOAD_CAP_OPENMP_400 ones first, and set
|
|
|
2554 NUM_DEVICES_OPENMP. */
|
|
|
2555 struct gomp_device_descr *devices_s
|
|
|
2556 = malloc (num_devices * sizeof (struct gomp_device_descr));
|
|
|
2557 if (!devices_s)
|
|
|
2558 {
|
|
|
2559 num_devices = 0;
|
|
|
2560 free (devices);
|
|
|
2561 devices = NULL;
|
|
|
2562 }
|
|
|
2563 num_devices_openmp = 0;
|
|
|
2564 for (i = 0; i < num_devices; i++)
|
|
|
2565 if (devices[i].capabilities & GOMP_OFFLOAD_CAP_OPENMP_400)
|
|
|
2566 devices_s[num_devices_openmp++] = devices[i];
|
|
|
2567 int num_devices_after_openmp = num_devices_openmp;
|
|
|
2568 for (i = 0; i < num_devices; i++)
|
|
|
2569 if (!(devices[i].capabilities & GOMP_OFFLOAD_CAP_OPENMP_400))
|
|
|
2570 devices_s[num_devices_after_openmp++] = devices[i];
|
|
|
2571 free (devices);
|
|
|
2572 devices = devices_s;
|
|
|
2573
|
|
|
2574 for (i = 0; i < num_devices; i++)
|
|
|
2575 {
|
|
|
2576 /* The 'devices' array can be moved (by the realloc call) until we have
|
|
|
2577 found all the plugins, so registering with the OpenACC runtime (which
|
|
|
2578 takes a copy of the pointer argument) must be delayed until now. */
|
|
|
2579 if (devices[i].capabilities & GOMP_OFFLOAD_CAP_OPENACC_200)
|
|
|
2580 goacc_register (&devices[i]);
|
|
|
2581 }
|
|
|
2582
|
|
|
2583 if (atexit (gomp_target_fini) != 0)
|
|
|
2584 gomp_fatal ("atexit failed");
|
|
|
2585 }
|
|
|
2586
|
|
|
2587 #else /* PLUGIN_SUPPORT */
|
|
|
2588 /* If dlfcn.h is unavailable we always fallback to host execution.
|
|
|
2589 GOMP_target* routines are just stubs for this case. */
|
|
|
2590 static void
|
|
|
2591 gomp_target_init (void)
|
|
|
2592 {
|
|
|
2593 }
|
|
|
2594 #endif /* PLUGIN_SUPPORT */
|
