From 094db6beb9cea0aedbde326f271d2b6fab762b1d Mon Sep 17 00:00:00 2001 From: Cesar Philippidis Date: Wed, 1 Aug 2018 07:09:56 -0700 Subject: [PATCH] [PATCH] Remove use of 'struct map' from plugin (nvptx) libgomp/ * plugin/plugin-nvptx.c (struct map): Removed. (map_init, map_pop): Remove use of struct map. (map_push): Likewise and change argument list. * testsuite/libgomp.oacc-c-c++-common/mapping-1.c: New Co-Authored-By: James Norris From-SVN: r263212 --- libgomp/ChangeLog | 8 +++ libgomp/plugin/plugin-nvptx.c | 33 ++-------- .../libgomp.oacc-c-c++-common/mapping-1.c | 63 +++++++++++++++++++ 3 files changed, 77 insertions(+), 27 deletions(-) create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/mapping-1.c diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog index 5ac69e85a99..cc5325fe8a8 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,3 +1,11 @@ +2018-08-01 Cesar Philippidis + James Norris + + * plugin/plugin-nvptx.c (struct map): Removed. + (map_init, map_pop): Remove use of struct map. (map_push): + Likewise and change argument list. + * testsuite/libgomp.oacc-c-c++-common/mapping-1.c: New + 2018-08-01 Tom de Vries * plugin/cuda-lib.def: New file. Factor out of ... diff --git a/libgomp/plugin/plugin-nvptx.c b/libgomp/plugin/plugin-nvptx.c index 83176ce07a0..cc465b4addb 100644 --- a/libgomp/plugin/plugin-nvptx.c +++ b/libgomp/plugin/plugin-nvptx.c @@ -177,13 +177,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) { @@ -217,16 +210,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); @@ -244,37 +233,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 == s->h_prev); + 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); @@ -1210,7 +1189,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 00000000000..593e7d4d553 --- /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); +} -- 2.30.2