From patchwork Tue Sep 23 18:17:36 2014 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Julian Brown X-Patchwork-Id: 392613 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Received: from sourceware.org (server1.sourceware.org [209.132.180.131]) (using TLSv1.2 with cipher ECDHE-RSA-AES256-GCM-SHA384 (256/256 bits)) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 52CA6140078 for ; Wed, 24 Sep 2014 04:18:01 +1000 (EST) DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:date :from:to:subject:message-id:mime-version:content-type; q=dns; s= default; b=WB4SYWsP1kzWhr79c7UtIp+aaMsHvRnlkWzGISgAhq0vXW7qFNnVC YjJ+pxs4q8DGwd05h9flq3EzyjYDnhVw5q9lgmHFdGXjkVKM7P5PQmUnXAep/ElG AWKfOg8HfboTnnamYxx5gjTiXdWeGycU61KidL+4Wkp4ZgMirder7s= DKIM-Signature: v=1; a=rsa-sha1; c=relaxed; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:date :from:to:subject:message-id:mime-version:content-type; s= default; bh=LvQqRYA2o1xoPayPYODHhlI99CI=; b=kWDUOXpPjQNr6vFWCU1Q fx/zZVZwMjxoDM7UCR+t0b7J8MVxexwpBbfTFaO561tPmAnnQTl8ka6jTYYHxI3w KMslVtL9QLguIzmJOw8jsmIjZ2dxoMrxQFKkpiZODD9ObDYA9zAsepOCYYeeWYxq Zv4+kM0Wg3qWKo/q3GKsWKA= Received: (qmail 12143 invoked by alias); 23 Sep 2014 18:17:53 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Unsubscribe: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Delivered-To: mailing list gcc-patches@gcc.gnu.org Received: (qmail 12124 invoked by uid 89); 23 Sep 2014 18:17:52 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-1.9 required=5.0 tests=AWL, BAYES_00 autolearn=ham version=3.3.2 X-HELO: relay1.mentorg.com Received: from relay1.mentorg.com (HELO relay1.mentorg.com) (192.94.38.131) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Tue, 23 Sep 2014 18:17:47 +0000 Received: from nat-ies.mentorg.com ([192.94.31.2] helo=SVR-IES-FEM-01.mgc.mentorg.com) by relay1.mentorg.com with esmtp id 1XWUem-0002ZO-2F from Julian_Brown@mentor.com for gcc-patches@gcc.gnu.org; Tue, 23 Sep 2014 11:17:44 -0700 Received: from octopus (137.202.0.76) by SVR-IES-FEM-01.mgc.mentorg.com (137.202.0.104) with Microsoft SMTP Server id 14.3.181.6; Tue, 23 Sep 2014 19:17:42 +0100 Date: Tue, 23 Sep 2014 19:17:36 +0100 From: Julian Brown To: Subject: [PATCH 1/10] OpenACC 2.0 support for libgomp - offloading support Message-ID: <20140923191736.5bdf18f1@octopus> MIME-Version: 1.0 X-IsSubscribed: yes This patch is by Jakub Jelinek, and was originally posted here: https://gcc.gnu.org/ml/gcc-patches/2013-09/msg01098.html Parts of the patch subsequently landed on mainline as part of the following patch: https://gcc.gnu.org/ml/gcc-patches/2013-10/msg00505.html But not the OpenMP "target" parts. This patch therefore contains the delta between those two patches. Julian xxxx-xx-xx Jakub Jelinek libgomp/ * splay-tree.h: New file. * target.c (splay_tree_node, splay_tree, splay_tree_key): New typedefs. (struct target_mem_desc, struct splay_tree_key_s): New structures. (splay_compare): New inline function. * libgomp.h (gomp_get_num_devices): Add prototype. (gomp_get_num_devices): Add FIXME comment. (resolve_device): Use default_device_var ICV. Add temporarily magic testing device number 257. (dev_splay_tree, dev_env_lock): New variables. (gomp_map_vars_existing, gomp_map_vars, gomp_unmap_tgt, gomp_unmap_vars, gomp_update): New functions. (GOMP_target, GOMP_target_data, GOMP_target_end_data, GOMP_target_update): Add support for magic testing device number 257. commit fc39aa98eba906466226c17fb455e57ebcfc1bc6 Author: Julian Brown Date: Fri Sep 19 08:33:05 2014 -0700 Delta between upstream and gomp-4_0-branch version of r202620: 2013-09-16 Jakub Jelinek * splay-tree.h: New file. * target.c: Include stdbool.h. (splay_tree_node, splay_tree, splay_tree_key): New typedefs. (struct target_mem_desc, struct splay_tree_key_s): New structures. (splay_compare): New inline function. (gomp_get_num_devices): New function. (resolve_device): Use default_device_var ICV. Add temporarily magic testing device number 257. (dev_splay_tree, dev_env_lock): New variables. (gomp_map_vars_existing, gomp_map_vars, gomp_unmap_tgt, gomp_unmap_vars, gomp_update): New functions. (GOMP_target, GOMP_target_data, GOMP_target_end_data, GOMP_target_update): Add support for magic testing device number 257. * libgomp.h (struct target_mem_desc): Forward declare. (struct gomp_task_icv): Add default_device_var and target_data. (gomp_get_num_devices): New prototype. * env.c (gomp_global_icv): Add default_device_var initializer. (parse_int): New function. (handle_omp_display_env): Print OMP_DEFAULT_DEVICE. (initialize_env): Initialize default_device_var. (omp_set_default_device): Set default_device_var ICV. (omp_get_default_device): Query default_device_var ICV. (omp_get_num_devices): Call gomp_get_num_devices. (omp_get_num_teams, omp_get_team_num, omp_is_initial_device): Add comments. diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h index a1482cc..d53a326 100644 --- a/libgomp/libgomp.h +++ b/libgomp/libgomp.h @@ -608,6 +608,10 @@ extern void gomp_free_thread (void *); extern int gomp_get_num_devices (void); +/* target.c */ + +extern int gomp_get_num_devices (void); + /* work.c */ extern void gomp_init_work_share (struct gomp_work_share *, bool, unsigned); diff --git a/libgomp/splay-tree.h b/libgomp/splay-tree.h new file mode 100644 index 0000000..04a71d1 --- /dev/null +++ b/libgomp/splay-tree.h @@ -0,0 +1,232 @@ +/* A splay-tree datatype. + Copyright 1998-2013 + Free Software Foundation, Inc. + Contributed by Mark Mitchell (mark@markmitchell.com). + + This file is part of the GNU OpenMP Library (libgomp). + + Libgomp is free software; you can redistribute it and/or modify it + under the terms of the GNU General Public License as published by + the Free Software Foundation; either version 3, or (at your option) + any later version. + + Libgomp is distributed in the hope that it will be useful, but WITHOUT ANY + WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS + FOR A PARTICULAR PURPOSE. See the GNU General Public License for + more details. + + Under Section 7 of GPL version 3, you are granted additional + permissions described in the GCC Runtime Library Exception, version + 3.1, as published by the Free Software Foundation. + + You should have received a copy of the GNU General Public License and + a copy of the GCC Runtime Library Exception along with this program; + see the files COPYING3 and COPYING.RUNTIME respectively. If not, see + . */ + +/* The splay tree code copied from include/splay-tree.h and adjusted, + so that all the data lives directly in splay_tree_node_s structure + and no extra allocations are needed. + + Files including this header should before including it add: +typedef struct splay_tree_node_s *splay_tree_node; +typedef struct splay_tree_s *splay_tree; +typedef struct splay_tree_key_s *splay_tree_key; + define splay_tree_key_s structure, and define + splay_compare inline function. */ + +/* For an easily readable description of splay-trees, see: + + Lewis, Harry R. and Denenberg, Larry. Data Structures and Their + Algorithms. Harper-Collins, Inc. 1991. + + The major feature of splay trees is that all basic tree operations + are amortized O(log n) time for a tree with n nodes. */ + +/* The nodes in the splay tree. */ +struct splay_tree_node_s { + struct splay_tree_key_s key; + /* The left and right children, respectively. */ + splay_tree_node left; + splay_tree_node right; +}; + +/* The splay tree. */ +struct splay_tree_s { + splay_tree_node root; +}; + +/* Rotate the edge joining the left child N with its parent P. PP is the + grandparents' pointer to P. */ + +static inline void +rotate_left (splay_tree_node *pp, splay_tree_node p, splay_tree_node n) +{ + splay_tree_node tmp; + tmp = n->right; + n->right = p; + p->left = tmp; + *pp = n; +} + +/* Rotate the edge joining the right child N with its parent P. PP is the + grandparents' pointer to P. */ + +static inline void +rotate_right (splay_tree_node *pp, splay_tree_node p, splay_tree_node n) +{ + splay_tree_node tmp; + tmp = n->left; + n->left = p; + p->right = tmp; + *pp = n; +} + +/* Bottom up splay of KEY. */ + +static void +splay_tree_splay (splay_tree sp, splay_tree_key key) +{ + if (sp->root == NULL) + return; + + do { + int cmp1, cmp2; + splay_tree_node n, c; + + n = sp->root; + cmp1 = splay_compare (key, &n->key); + + /* Found. */ + if (cmp1 == 0) + return; + + /* Left or right? If no child, then we're done. */ + if (cmp1 < 0) + c = n->left; + else + c = n->right; + if (!c) + return; + + /* Next one left or right? If found or no child, we're done + after one rotation. */ + cmp2 = splay_compare (key, &c->key); + if (cmp2 == 0 + || (cmp2 < 0 && !c->left) + || (cmp2 > 0 && !c->right)) + { + if (cmp1 < 0) + rotate_left (&sp->root, n, c); + else + rotate_right (&sp->root, n, c); + return; + } + + /* Now we have the four cases of double-rotation. */ + if (cmp1 < 0 && cmp2 < 0) + { + rotate_left (&n->left, c, c->left); + rotate_left (&sp->root, n, n->left); + } + else if (cmp1 > 0 && cmp2 > 0) + { + rotate_right (&n->right, c, c->right); + rotate_right (&sp->root, n, n->right); + } + else if (cmp1 < 0 && cmp2 > 0) + { + rotate_right (&n->left, c, c->right); + rotate_left (&sp->root, n, n->left); + } + else if (cmp1 > 0 && cmp2 < 0) + { + rotate_left (&n->right, c, c->left); + rotate_right (&sp->root, n, n->right); + } + } while (1); +} + +/* Insert a new NODE into SP. The NODE shouldn't exist in the tree. */ + +static void +splay_tree_insert (splay_tree sp, splay_tree_node node) +{ + int comparison = 0; + + splay_tree_splay (sp, &node->key); + + if (sp->root) + comparison = splay_compare (&sp->root->key, &node->key); + + if (sp->root && comparison == 0) + abort (); + else + { + /* Insert it at the root. */ + if (sp->root == NULL) + node->left = node->right = NULL; + else if (comparison < 0) + { + node->left = sp->root; + node->right = node->left->right; + node->left->right = NULL; + } + else + { + node->right = sp->root; + node->left = node->right->left; + node->right->left = NULL; + } + + sp->root = node; + } +} + +/* Remove node with KEY from SP. It is not an error if it did not exist. */ + +static void +splay_tree_remove (splay_tree sp, splay_tree_key key) +{ + splay_tree_splay (sp, key); + + if (sp->root && splay_compare (&sp->root->key, key) == 0) + { + splay_tree_node left, right; + + left = sp->root->left; + right = sp->root->right; + + /* One of the children is now the root. Doesn't matter much + which, so long as we preserve the properties of the tree. */ + if (left) + { + sp->root = left; + + /* If there was a right child as well, hang it off the + right-most leaf of the left child. */ + if (right) + { + while (left->right) + left = left->right; + left->right = right; + } + } + else + sp->root = right; + } +} + +/* Lookup KEY in SP, returning NODE if present, and NULL + otherwise. */ + +static splay_tree_key +splay_tree_lookup (splay_tree sp, splay_tree_key key) +{ + splay_tree_splay (sp, key); + + if (sp->root && splay_compare (&sp->root->key, key) == 0) + return &sp->root->key; + else + return NULL; +} diff --git a/libgomp/target.c b/libgomp/target.c index 46acc58..9b759b1 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -31,12 +31,379 @@ #include #include +/* Forward declaration for a node in the tree. */ +typedef struct splay_tree_node_s *splay_tree_node; +typedef struct splay_tree_s *splay_tree; +typedef struct splay_tree_key_s *splay_tree_key; + +struct target_mem_desc { + /* Reference count. */ + uintptr_t refcount; + /* All the splay nodes allocated together. */ + splay_tree_node array; + /* Start of the target region. */ + uintptr_t tgt_start; + /* End of the targer region. */ + uintptr_t tgt_end; + /* Handle to free. */ + void *to_free; + /* Previous target_mem_desc. */ + struct target_mem_desc *prev; + /* Number of items in following list. */ + size_t list_count; + /* List of splay keys to remove (or decrease refcount) + at the end of region. */ + splay_tree_key list[]; +}; + +struct splay_tree_key_s { + /* Address of the host object. */ + uintptr_t host_start; + /* Address immediately after the host object. */ + uintptr_t host_end; + /* Descriptor of the target memory. */ + struct target_mem_desc *tgt; + /* Offset from tgt->tgt_start to the start of the target object. */ + uintptr_t tgt_offset; + /* Reference count. */ + uintptr_t refcount; + /* True if data should be copied from device to host at the end. */ + bool copy_from; +}; + +/* The comparison function. */ + +static int +splay_compare (splay_tree_key x, splay_tree_key y) +{ + if (x->host_start == x->host_end + && y->host_start == y->host_end) + return 0; + if (x->host_end <= y->host_start) + return -1; + if (x->host_start >= y->host_end) + return 1; + return 0; +} + +#include "splay-tree.h" + attribute_hidden int gomp_get_num_devices (void) { + /* FIXME: Scan supported accelerators when called the first time. */ return 0; } +static int +resolve_device (int device) +{ + if (device == -1) + { + struct gomp_task_icv *icv = gomp_icv (false); + device = icv->default_device_var; + } + /* FIXME: Temporary hack for testing non-shared address spaces on host. */ + if (device == 257) + return 257; + if (device >= gomp_get_num_devices ()) + return -1; + return -1; +} + +/* These variables would be per-accelerator (which doesn't have shared address + space. */ +static struct splay_tree_s dev_splay_tree; +static gomp_mutex_t dev_env_lock; + +/* Handle the case where splay_tree_lookup found oldn for newn. + Helper function of gomp_map_vars. */ + +static inline void +gomp_map_vars_existing (splay_tree_key oldn, splay_tree_key newn, + unsigned char kind) +{ + if (oldn->host_start > newn->host_start + || oldn->host_end < newn->host_end) + gomp_fatal ("Trying to map into device [%p..%p) object when" + "[%p..%p) is already mapped", + (void *) newn->host_start, (void *) newn->host_end, + (void *) oldn->host_start, (void *) oldn->host_end); + if (((kind & 7) == 2 || (kind & 7) == 3) + && !oldn->copy_from + && oldn->host_start == newn->host_start + && oldn->host_end == newn->host_end) + oldn->copy_from = true; + oldn->refcount++; +} + +static struct target_mem_desc * +gomp_map_vars (size_t mapnum, void **hostaddrs, size_t *sizes, + unsigned char *kinds, bool is_target) +{ + size_t i, tgt_align, tgt_size, not_found_cnt = 0; + struct splay_tree_key_s cur_node; + struct target_mem_desc *tgt + = gomp_malloc (sizeof (*tgt) + sizeof (tgt->list[0]) * mapnum); + tgt->list_count = mapnum; + tgt->refcount = 1; + + if (mapnum == 0) + return tgt; + + tgt_align = sizeof (void *); + tgt_size = 0; + if (is_target) + { + size_t align = 4 * sizeof (void *); + tgt_align = align; + tgt_size = mapnum * sizeof (void *); + } + gomp_mutex_lock (&dev_env_lock); + for (i = 0; i < mapnum; i++) + { + cur_node.host_start = (uintptr_t) hostaddrs[i]; + if ((kinds[i] & 7) != 4) + cur_node.host_end = cur_node.host_start + sizes[i]; + else + cur_node.host_end = cur_node.host_start + sizeof (void *); + splay_tree_key n = splay_tree_lookup (&dev_splay_tree, &cur_node); + if (n) + { + tgt->list[i] = n; + gomp_map_vars_existing (n, &cur_node, kinds[i]); + } + else + { + size_t align = (size_t) 1 << (kinds[i] >> 3); + tgt->list[i] = NULL; + not_found_cnt++; + if (tgt_align < align) + tgt_align = align; + tgt_size = (tgt_size + align - 1) & ~(align - 1); + tgt_size += cur_node.host_end - cur_node.host_start; + } + } + + if (not_found_cnt || is_target) + { + /* FIXME: This would be accelerator memory allocation, not + host, and should allocate tgt_align aligned tgt_size block + of memory. */ + tgt->to_free = gomp_malloc (tgt_size + tgt_align - 1); + tgt->tgt_start = (uintptr_t) tgt->to_free; + tgt->tgt_start = (tgt->tgt_start + tgt_align - 1) & ~(tgt_align - 1); + tgt->tgt_end = tgt->tgt_start + tgt_size; + } + + tgt_size = 0; + if (is_target) + tgt_size = mapnum * sizeof (void *); + + if (not_found_cnt) + { + tgt->array = gomp_malloc (not_found_cnt * sizeof (*tgt->array)); + splay_tree_node array = tgt->array; + + for (i = 0; i < mapnum; i++) + if (tgt->list[i] == NULL) + { + splay_tree_key k = &array->key; + k->host_start = (uintptr_t) hostaddrs[i]; + if ((kinds[i] & 7) != 4) + k->host_end = k->host_start + sizes[i]; + else + k->host_end = k->host_start + sizeof (void *); + splay_tree_key n + = splay_tree_lookup (&dev_splay_tree, k); + if (n) + { + tgt->list[i] = n; + gomp_map_vars_existing (n, k, kinds[i]); + } + else + { + size_t align = (size_t) 1 << (kinds[i] >> 3); + tgt->list[i] = k; + tgt_size = (tgt_size + align - 1) & ~(align - 1); + k->tgt = tgt; + k->tgt_offset = tgt_size; + tgt_size += k->host_end - k->host_start; + if ((kinds[i] & 7) == 2 || (kinds[i] & 7) == 3) + k->copy_from = true; + k->refcount = 1; + tgt->refcount++; + array->left = NULL; + array->right = NULL; + splay_tree_insert (&dev_splay_tree, array); + switch (kinds[i] & 7) + { + case 0: /* ALLOC */ + case 2: /* FROM */ + break; + case 1: /* TO */ + case 3: /* TOFROM */ + /* FIXME: This is supposed to be copy from host to device + memory. Perhaps add some smarts, like if copying + several adjacent fields from host to target, use some + host buffer to avoid sending each var individually. */ + memcpy ((void *) (tgt->tgt_start + k->tgt_offset), + (void *) k->host_start, + k->host_end - k->host_start); + break; + case 4: /* POINTER */ + cur_node.host_start + = (uintptr_t) *(void **) k->host_start; + /* Add bias to the pointer value. */ + cur_node.host_start += sizes[i]; + cur_node.host_end = cur_node.host_start + 1; + n = splay_tree_lookup (&dev_splay_tree, &cur_node); + if (n == NULL) + { + /* Could be possibly zero size array section. */ + cur_node.host_end--; + n = splay_tree_lookup (&dev_splay_tree, &cur_node); + if (n == NULL) + { + cur_node.host_start--; + n = splay_tree_lookup (&dev_splay_tree, &cur_node); + cur_node.host_start++; + } + } + if (n == NULL) + gomp_fatal ("Pointer target of array section " + "wasn't mapped"); + cur_node.host_start -= n->host_start; + cur_node.tgt_offset = n->tgt->tgt_start + n->tgt_offset + + cur_node.host_start; + /* At this point tgt_offset is target address of the + array section. Now subtract bias to get what we want + to initialize the pointer with. */ + cur_node.tgt_offset -= sizes[i]; + /* FIXME: host to device copy, see above FIXME comment. */ + memcpy ((void *) (tgt->tgt_start + k->tgt_offset), + (void *) &cur_node.tgt_offset, + sizeof (void *)); + break; + } + array++; + } + } + } + if (is_target) + { + for (i = 0; i < mapnum; i++) + { + cur_node.tgt_offset = tgt->list[i]->tgt->tgt_start + + tgt->list[i]->tgt_offset; + /* FIXME: host to device copy, see above FIXME comment. */ + memcpy ((void *) (tgt->tgt_start + i * sizeof (void *)), + (void *) &cur_node.tgt_offset, + sizeof (void *)); + } + } + + gomp_mutex_unlock (&dev_env_lock); + return tgt; +} + +static void +gomp_unmap_tgt (struct target_mem_desc *tgt) +{ + /* FIXME: Deallocate on target the tgt->tgt_start .. tgt->tgt_end + region. */ + if (tgt->tgt_end) + free (tgt->to_free); + + free (tgt->array); + free (tgt); +} + +static void +gomp_unmap_vars (struct target_mem_desc *tgt) +{ + if (tgt->list_count == 0) + { + free (tgt); + return; + } + + size_t i; + gomp_mutex_lock (&dev_env_lock); + for (i = 0; i < tgt->list_count; i++) + if (tgt->list[i]->refcount > 1) + tgt->list[i]->refcount--; + else + { + splay_tree_key k = tgt->list[i]; + if (k->copy_from) + /* FIXME: device to host copy. */ + memcpy ((void *) k->host_start, + (void *) (k->tgt->tgt_start + k->tgt_offset), + k->host_end - k->host_start); + splay_tree_remove (&dev_splay_tree, k); + if (k->tgt->refcount > 1) + k->tgt->refcount--; + else + gomp_unmap_tgt (k->tgt); + } + + if (tgt->refcount > 1) + tgt->refcount--; + else + gomp_unmap_tgt (tgt); + gomp_mutex_unlock (&dev_env_lock); +} + +static void +gomp_update (size_t mapnum, void **hostaddrs, size_t *sizes, + unsigned char *kinds) +{ + size_t i; + struct splay_tree_key_s cur_node; + + if (mapnum == 0) + return; + + gomp_mutex_lock (&dev_env_lock); + for (i = 0; i < mapnum; i++) + if (sizes[i]) + { + cur_node.host_start = (uintptr_t) hostaddrs[i]; + cur_node.host_end = cur_node.host_start + sizes[i]; + splay_tree_key n = splay_tree_lookup (&dev_splay_tree, &cur_node); + if (n) + { + if (n->host_start > cur_node.host_start + || n->host_end < cur_node.host_end) + gomp_fatal ("Trying to update [%p..%p) object when" + "only [%p..%p) is mapped", + (void *) cur_node.host_start, + (void *) cur_node.host_end, + (void *) n->host_start, + (void *) n->host_end); + if ((kinds[i] & 7) == 1) + /* FIXME: host to device copy. */ + memcpy ((void *) (n->tgt->tgt_start + n->tgt_offset + + cur_node.host_start - n->host_start), + (void *) cur_node.host_start, + cur_node.host_end - cur_node.host_start); + else if ((kinds[i] & 7) == 2) + /* FIXME: device to host copy. */ + memcpy ((void *) cur_node.host_start, + (void *) (n->tgt->tgt_start + n->tgt_offset + + cur_node.host_start - n->host_start), + cur_node.host_end - cur_node.host_start); + } + else + gomp_fatal ("Trying to update [%p..%p) object that is not mapped", + (void *) cur_node.host_start, + (void *) cur_node.host_end); + } + gomp_mutex_unlock (&dev_env_lock); +} + + /* Called when encountering a target directive. If DEVICE is -1, it means use device-var ICV. If it is -2 (or any other value larger than last available hw device, use host fallback. @@ -52,35 +419,87 @@ GOMP_target (int device, void (*fn) (void *), const void *openmp_target, size_t mapnum, void **hostaddrs, size_t *sizes, unsigned char *kinds) { - /* Host fallback. */ - struct gomp_thread old_thr, *thr = gomp_thread (); - old_thr = *thr; - memset (thr, '\0', sizeof (*thr)); - if (gomp_places_list) + device = resolve_device (device); + if (device == -1) + { + /* Host fallback. */ + struct gomp_thread old_thr, *thr = gomp_thread (); + old_thr = *thr; + memset (thr, '\0', sizeof (*thr)); + if (gomp_places_list) + { + thr->place = old_thr.place; + thr->ts.place_partition_len = gomp_places_list_len; + } + fn (hostaddrs); + gomp_free_thread (thr); + *thr = old_thr; + return; + } + if (device == 257) { - thr->place = old_thr.place; - thr->ts.place_partition_len = gomp_places_list_len; + struct target_mem_desc *tgt + = gomp_map_vars (mapnum, hostaddrs, sizes, kinds, true); + fn ((void *) tgt->tgt_start); + gomp_unmap_vars (tgt); } - fn (hostaddrs); - gomp_free_thread (thr); - *thr = old_thr; } void GOMP_target_data (int device, const void *openmp_target, size_t mapnum, void **hostaddrs, size_t *sizes, unsigned char *kinds) { + device = resolve_device (device); + if (device == -1) + { + /* Host fallback. */ + struct gomp_task_icv *icv = gomp_icv (false); + if (icv->target_data) + { + /* Even when doing a host fallback, if there are any active + #pragma omp target data constructs, need to remember the + new #pragma omp target data, otherwise GOMP_target_end_data + would get out of sync. */ + struct target_mem_desc *tgt + = gomp_map_vars (0, NULL, NULL, NULL, false); + tgt->prev = icv->target_data; + icv->target_data = tgt; + } + return; + } + + if (device == 257) + { + struct target_mem_desc *tgt + = gomp_map_vars (mapnum, hostaddrs, sizes, kinds, false); + struct gomp_task_icv *icv = gomp_icv (true); + tgt->prev = icv->target_data; + icv->target_data = tgt; + } } void GOMP_target_end_data (void) { + struct gomp_task_icv *icv = gomp_icv (false); + if (icv->target_data) + { + struct target_mem_desc *tgt = icv->target_data; + icv->target_data = tgt->prev; + gomp_unmap_vars (tgt); + } } void GOMP_target_update (int device, const void *openmp_target, size_t mapnum, void **hostaddrs, size_t *sizes, unsigned char *kinds) { + device = resolve_device (device); + if (device == -1) + return; + + if (device == 257) + gomp_update (mapnum, hostaddrs, sizes, kinds); } void