[PATCH,nvptx] Remove use of 'struct map' from plugin (nvptx)

[Thomas Schwinge]

Hi!

On Tue, 31 Jul 2018 08:12:51 -0700, Cesar Philippidis <cesar_philippidis@mentor.com> wrote:
> This is an old patch which removes the struct map from the nvptx plugin.

(This got committed to trunk in r263212.)

> I believe at one point this was supposed to be used to manage async data
> mappings, but in practice that never worked out.

The original submission seems to be
<http://mid.mail-archive.com/566EC830.8030508@codesourcery.com>:

| The attached patch fixes an issue in the managing of
| the page-locked buffer which holds the kernel launch
| mappings. In the process of fixing the issue, I discovered
| that 'struct map' was no longer needed, so that has
| been removed as well.

I can't tell/remember what the "issue in the managing of the page-locked
buffer which holds the kernel launch mappings" would've been.

> --- /dev/null
> +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/mapping-1.c

I wanted to note that on all branches down to GCC 5 this test case PASSes
already without any libgomp code changes.

> @@ -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;

The rationale for the number 48 doesn't seem right -- the first
'parallel' construct has four mappings, so four times eight bytes is 32
bytes on x86_64?  Ha, or maybe this applies to the 'struct map', so it's
these four plus the 'int async', 'size_t size'?

Anyway, I can't tell what this is trying to achieve?

> +
> +  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;
> +	    }
> +	}
> +    }

Maybe filling up some improperly-managed statically-sized data structure
(said "page-locked buffer"?), which then...

> +
> +#pragma acc parallel copy (a, b, c, d, e, f)

... would overflow here?

> +  {
> +    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);
> +}

Anyway -- the libgomp code has been cleaned up; the test case seems like
it can be disregarded.


Grüße
 Thomas
-------------- next part --------------
A non-text attachment was scrubbed...
Name: signature.asc
Type: application/pgp-signature
Size: 658 bytes
Desc: not available
URL: <http://gcc.gnu.org/pipermail/gcc-patches/attachments/20190506/5c9e6c1c/attachment.sig>