From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: (qmail 84546 invoked by alias); 12 Jul 2016 22:33:13 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Received: (qmail 81344 invoked by uid 89); 12 Jul 2016 22:33:08 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-1.7 required=5.0 tests=AWL,BAYES_00,RCVD_IN_DNSWL_NONE,SPF_PASS autolearn=ham version=3.3.2 spammy=201512, norris, Norris 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 (AES256-GCM-SHA384 encrypted) ESMTPS; Tue, 12 Jul 2016 22:32:57 +0000 Received: from svr-orw-fem-02x.mgc.mentorg.com ([147.34.96.206] helo=SVR-ORW-FEM-02.mgc.mentorg.com) by relay1.mentorg.com with esmtp id 1bN6EX-0007Tz-Rc from Cesar_Philippidis@mentor.com ; Tue, 12 Jul 2016 15:32:53 -0700 Received: from [127.0.0.1] (147.34.91.1) by svr-orw-fem-02.mgc.mentorg.com (147.34.96.168) with Microsoft SMTP Server id 14.3.224.2; Tue, 12 Jul 2016 15:32:53 -0700 From: Cesar Philippidis Subject: [PATCH] Remove struct map from plugin-nvptx To: "gcc-patches@gcc.gnu.org" , Jakub Jelinek CC: james norris Message-ID: <57857015.8020801@codesourcery.com> Date: Tue, 12 Jul 2016 22:33:00 -0000 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:38.0) Gecko/20100101 Thunderbird/38.8.0 MIME-Version: 1.0 Content-Type: multipart/mixed; boundary="------------040009090108040807050603" X-SW-Source: 2016-07/txt/msg00690.txt.bz2 --------------040009090108040807050603 Content-Type: text/plain; charset="utf-8" Content-Transfer-Encoding: 7bit Content-length: 669 I'm not sure why the nvptx offloading plugin introduced struct map in the first place, but its current usage is both unnecessary and buggy. For instance, it doesn't properly free data mappings, and that results in bogus duplicate data mapping errors. This patch removes that struct altogether. Jim originally posted this patch back in December, but it looks like it never made it's way into trunk. This patch is in gomp4 though, but I had to rebase it to trunk because trunk contains some runtime async changes that aren't in gomp4 yet. Here's the posting to the original patch . Is this OK for trunk? Cesar --------------040009090108040807050603 Content-Type: text/x-patch; name="libgomp-map-removal.diff" Content-Transfer-Encoding: 7bit Content-Disposition: attachment; filename="libgomp-map-removal.diff" Content-length: 3845 2016-07-12 Cesar Philippidis James Norris libgomp/ * plugin/plugin-nvptx.c (struct map): Delete. (map_pop): Remove use of struct map. (map_push): Likewise. Remove async argument. (nvptx_exec): Update call to map_push. * testsuite/libgomp.oacc-c-c++-common/mapping-1.c: New test. diff --git a/libgomp/plugin/plugin-nvptx.c b/libgomp/plugin/plugin-nvptx.c index 327500c..7245e67 100644 --- a/libgomp/plugin/plugin-nvptx.c +++ b/libgomp/plugin/plugin-nvptx.c @@ -119,13 +119,6 @@ struct nvptx_thread struct ptx_device *ptx_dev; }; -struct map -{ - int async; - size_t size; - char mappings[0]; -}; - static bool map_init (struct ptx_stream *s) { @@ -159,16 +152,12 @@ map_fini (struct ptx_stream *s) static void map_pop (struct ptx_stream *s) { - struct map *m; - assert (s != NULL); assert (s->h_next); assert (s->h_prev); assert (s->h_tail); - m = s->h_tail; - - s->h_tail += m->size; + s->h_tail = s->h_next; if (s->h_tail >= s->h_end) s->h_tail = s->h_begin + (int) (s->h_tail - s->h_end); @@ -186,37 +175,27 @@ map_pop (struct ptx_stream *s) } static void -map_push (struct ptx_stream *s, int async, size_t size, void **h, void **d) +map_push (struct ptx_stream *s, size_t size, void **h, void **d) { int left; int offset; - struct map *m; assert (s != NULL); left = s->h_end - s->h_next; - size += sizeof (struct map); assert (s->h_prev); assert (s->h_next); if (size >= left) { - m = s->h_prev; - m->size += left; - s->h_next = s->h_begin; - - if (s->h_next + size > s->h_end) - GOMP_PLUGIN_fatal ("unable to push map"); + assert (s->h_next + size > s->h_end); + s->h_next = s->h_prev = s->h_tail = s->h_begin; } assert (s->h_next); - m = s->h_next; - m->async = async; - m->size = size; - - offset = (void *)&m->mappings[0] - s->h; + offset = s->h_next - s->h; *d = (void *)(s->d + offset); *h = (void *)(s->h + offset); @@ -940,7 +919,7 @@ nvptx_exec (void (*fn), size_t mapnum, void **hostaddrs, void **devaddrs, /* This reserves a chunk of a pre-allocated page of memory mapped on both the host and the device. HP is a host pointer to the new chunk, and DP is the corresponding device pointer. */ - map_push (dev_str, async, mapnum * sizeof (void *), &hp, &dp); + map_push (dev_str, mapnum * sizeof (void *), &hp, &dp); GOMP_PLUGIN_debug (0, " %s: prepare mappings\n", __FUNCTION__); diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/mapping-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/mapping-1.c new file mode 100644 index 0000000..593e7d4 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/mapping-1.c @@ -0,0 +1,63 @@ +/* { dg-do run } */ + +#include +#include +#include + +/* Exercise the kernel launch argument mapping. */ + +int +main (int argc, char **argv) +{ + int a[256], b[256], c[256], d[256], e[256], f[256]; + int i; + int n; + + /* 48 is the size of the mappings for the first parallel construct. */ + n = sysconf (_SC_PAGESIZE) / 48 - 1; + + i = 0; + + for (i = 0; i < n; i++) + { + #pragma acc parallel copy (a, b, c, d) + { + int j; + + for (j = 0; j < 256; j++) + { + a[j] = j; + b[j] = j; + c[j] = j; + d[j] = j; + } + } + } + +#pragma acc parallel copy (a, b, c, d, e, f) + { + int j; + + for (j = 0; j < 256; j++) + { + a[j] = j; + b[j] = j; + c[j] = j; + d[j] = j; + e[j] = j; + f[j] = j; + } + } + + for (i = 0; i < 256; i++) + { + if (a[i] != i) abort(); + if (b[i] != i) abort(); + if (c[i] != i) abort(); + if (d[i] != i) abort(); + if (e[i] != i) abort(); + if (f[i] != i) abort(); + } + + exit (0); +} --------------040009090108040807050603--