[PATCH] Remove struct map from plugin-nvptx

[PATCH] Remove struct map from plugin-nvptx

[Cesar Philippidis]

[-- Attachment #1: Type: text/plain, Size: 669 bytes --]

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
<https://gcc.gnu.org/ml/gcc-patches/2015-12/msg01658.html>.

Is this OK for trunk?

Cesar

[-- Attachment #2: libgomp-map-removal.diff --]
[-- Type: text/x-patch, Size: 3845 bytes --]

2016-07-12  Cesar Philippidis  <cesar@codesourcery.com>
	    James Norris  <jnorris@codesourcery.com>

	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 <stdio.h>
+#include <stdlib.h>
+#include <unistd.h>
+
+/* 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);
+}