From patchwork Tue Oct 28 16:52:22 2014 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Julian Brown X-Patchwork-Id: 404299 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 1E3C014003E for ; Wed, 29 Oct 2014 03:52:43 +1100 (AEDT) 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:cc:subject:message-id:mime-version:content-type; q=dns; s=default; b=jL66uC6d+TBXxWjBtq5UjLsOjwLfI4wfDpYleK9GOQbxKAEohB riyk2r2d9rofqLfUqGhz/j852TlsHcFqLDhGHycFhxSzLeS+nFc2M6tjs9zJE810 CEXw3uNmVIeNS34ciOgHQKkm7KVThJCvSQPSGIn7tg4lYvVjlhHcuBrQI= 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:cc:subject:message-id:mime-version:content-type; s= default; bh=YoQesygG3CQUSlqIDvnxyr6L0U0=; b=k5m9RJJ8YGHdcIdYIxK0 zkogus/EeOXaikul+NDU2dDp5NvODKJWodzW4IoaE4ZoNQvavUny4iO8BZwU892v g8ejLR44FHGxAYn29jOfD/CT9PomkgHScQxEaLqPhZsfQ0ToXDQHHIl06hXP2EN5 ZAvBoVuBLMiPF7NuvRECIEM= Received: (qmail 30877 invoked by alias); 28 Oct 2014 16:52:36 -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 30868 invoked by uid 89); 28 Oct 2014 16:52:35 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-0.4 required=5.0 tests=AWL, BAYES_00, RCVD_IN_DNSWL_NONE, UNWANTED_LANGUAGE_BODY 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, 28 Oct 2014 16:52:34 +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 1XjA0U-0001Zp-38 from Julian_Brown@mentor.com ; Tue, 28 Oct 2014 09:52:30 -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, 28 Oct 2014 16:52:28 +0000 Date: Tue, 28 Oct 2014 16:52:22 +0000 From: Julian Brown To: CC: Thomas Schwinge , Jakub Jelinek Subject: [gomp4] Remove gomp_map_vars mem_map argument Message-ID: <20141028165222.2f9a415a@octopus> MIME-Version: 1.0 X-IsSubscribed: yes Hi, This patch removes the now-redundant gomp_memory_mapping argument from gomp_map_vars, introduced when OpenACC kept the structure in question in a different place from OpenMP. Both now keep the memory map in the gomp_device_descr structure, so there's no need to pass both that and the memory map to the function explicitly. OK for gomp4 branch? Thanks, Julian ChangeLog libgomp/ * target.c (gomp_map_vars): Remove MM argument. (GOMP_target, GOMP_target_data): Update calls to gomp_map_vars. * oacc-mem.c (acc_map_data, present_create_copy): Update calls to gomp_map_vars. * oacc-parallel.c (GOACC_parallel, GOACC_data_start): Likewise. * target.h (gomp_map_vars): Update prototype. commit 3afc4e592a6d8a796ec0c44bb8dc808b1392fd29 Author: Julian Brown Date: Tue Oct 28 09:17:01 2014 -0700 Remove gomp_map_vars mem_map argument diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c index d812f72..582a1e0 100644 --- a/libgomp/oacc-mem.c +++ b/libgomp/oacc-mem.c @@ -257,7 +257,7 @@ acc_map_data (void *h, void *d, size_t s) if (d != h) gomp_fatal ("cannot map data on shared-memory system"); - tgt = gomp_map_vars (NULL, NULL, 0, NULL, NULL, NULL, NULL, true, false); + tgt = gomp_map_vars (NULL, 0, NULL, NULL, NULL, NULL, true, false); } else { @@ -275,9 +275,8 @@ acc_map_data (void *h, void *d, size_t s) gomp_fatal ("device address [%p, +%d] is already mapped", (void *)d, (int)s); - tgt = gomp_map_vars ((struct gomp_device_descr *) acc_dev, - &acc_dev->mem_map, mapnum, &hostaddrs, - &devaddrs, &sizes, &kinds, true, false); + tgt = gomp_map_vars (acc_dev, mapnum, &hostaddrs, &devaddrs, &sizes, + &kinds, true, false); } tgt->prev = acc_dev->openacc.data_environ; @@ -383,9 +382,8 @@ present_create_copy (unsigned f, void *h, size_t s) else kinds = GOMP_MAP_ALLOC; - tgt = gomp_map_vars ((struct gomp_device_descr *) acc_dev, - &acc_dev->mem_map, mapnum, &hostaddrs, - NULL, &s, &kinds, true, false); + tgt = gomp_map_vars (acc_dev, mapnum, &hostaddrs, NULL, &s, &kinds, true, + false); gomp_mutex_lock (&acc_dev->mem_map.lock); diff --git a/libgomp/oacc-parallel.c b/libgomp/oacc-parallel.c index b787df7..1639244 100644 --- a/libgomp/oacc-parallel.c +++ b/libgomp/oacc-parallel.c @@ -173,9 +173,8 @@ GOACC_parallel (int device, void (*fn) (void *), const void *openmp_target, else tgt_fn = (void (*)) fn; - tgt = gomp_map_vars ((struct gomp_device_descr *) acc_dev, - &acc_dev->mem_map, mapnum, hostaddrs, - NULL, sizes, kinds, true, false); + tgt = gomp_map_vars (acc_dev, mapnum, hostaddrs, NULL, sizes, kinds, true, + false); devaddrs = alloca (sizeof (void *) * mapnum); for (i = 0; i < mapnum; i++) @@ -217,7 +216,7 @@ GOACC_data_start (int device, const void *openmp_target, size_t mapnum, if ((acc_dev->capabilities & TARGET_CAP_SHARED_MEM) || !if_clause_condition_value) { - tgt = gomp_map_vars (NULL, NULL, 0, NULL, NULL, NULL, NULL, true, false); + tgt = gomp_map_vars (NULL, 0, NULL, NULL, NULL, NULL, true, false); tgt->prev = thr->mapped_data; thr->mapped_data = tgt; @@ -225,9 +224,8 @@ GOACC_data_start (int device, const void *openmp_target, size_t mapnum, } gomp_notify (" %s: prepare mappings\n", __FUNCTION__); - tgt = gomp_map_vars ((struct gomp_device_descr *) acc_dev, - &acc_dev->mem_map, mapnum, hostaddrs, - NULL, sizes, kinds, true, false); + tgt = gomp_map_vars (acc_dev, mapnum, hostaddrs, NULL, sizes, kinds, true, + false); gomp_notify (" %s: mappings prepared\n", __FUNCTION__); tgt->prev = thr->mapped_data; thr->mapped_data = tgt; diff --git a/libgomp/target.c b/libgomp/target.c index 615ba6b..507488e 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -134,14 +134,14 @@ get_kind (bool is_openacc, void *kinds, int idx) } attribute_hidden struct target_mem_desc * -gomp_map_vars (struct gomp_device_descr *devicep, - struct gomp_memory_mapping *mm, size_t mapnum, - void **hostaddrs, void **devaddrs, size_t *sizes, - void *kinds, bool is_openacc, bool is_target) +gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum, + void **hostaddrs, void **devaddrs, size_t *sizes, void *kinds, + bool is_openacc, bool is_target) { size_t i, tgt_align, tgt_size, not_found_cnt = 0; const int rshift = is_openacc ? 8 : 3; const int typemask = is_openacc ? 0xff : 0x7; + struct gomp_memory_mapping *mm = &devicep->mem_map; struct splay_tree_key_s cur_node; struct target_mem_desc *tgt = gomp_malloc (sizeof (*tgt) + sizeof (tgt->list[0]) * mapnum); @@ -861,8 +861,8 @@ GOMP_target (int device, void (*fn) (void *), const void *openmp_target, gomp_mutex_unlock (&mm->lock); struct target_mem_desc *tgt_vars - = gomp_map_vars (devicep, &devicep->mem_map, mapnum, hostaddrs, NULL, - sizes, kinds, false, true); + = gomp_map_vars (devicep, mapnum, hostaddrs, NULL, sizes, kinds, false, + true); struct gomp_thread old_thr, *thr = gomp_thread (); old_thr = *thr; memset (thr, '\0', sizeof (*thr)); @@ -901,17 +901,15 @@ GOMP_target_data (int device, const void *openmp_target, size_t mapnum, new #pragma omp target data, otherwise GOMP_target_end_data would get out of sync. */ struct target_mem_desc *tgt - = gomp_map_vars (NULL, NULL, 0, NULL, NULL, NULL, NULL, false, - false); + = gomp_map_vars (NULL, 0, NULL, NULL, NULL, NULL, false, false); tgt->prev = icv->target_data; icv->target_data = tgt; } return; } - struct target_mem_desc *tgt - = gomp_map_vars (devicep, &devicep->mem_map, mapnum, hostaddrs, NULL, sizes, - kinds, false, false); + struct target_mem_desc *tgt = gomp_map_vars (devicep, mapnum, hostaddrs, NULL, + sizes, kinds, false, false); struct gomp_task_icv *icv = gomp_icv (true); tgt->prev = icv->target_data; icv->target_data = tgt; diff --git a/libgomp/target.h b/libgomp/target.h index b5773e2..d4c1120 100644 --- a/libgomp/target.h +++ b/libgomp/target.h @@ -199,10 +199,9 @@ struct gomp_device_descr }; extern struct target_mem_desc * -gomp_map_vars (struct gomp_device_descr *devicep, - struct gomp_memory_mapping *mm, size_t mapnum, - void **hostaddrs, void **devaddrs, size_t *sizes, - void *kinds, bool is_openacc, bool is_target); +gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum, + void **hostaddrs, void **devaddrs, size_t *sizes, void *kinds, + bool is_openacc, bool is_target); extern void gomp_copy_from_async (struct target_mem_desc *tgt);