[PR90247] Reconsider OpenACC implicit data attributes for pointers (was: [og7] Enable 0-length array data mappings for implicit data clauses)

[Thomas Schwinge]

Hi!

On Wed, 11 Oct 2017 07:31:04 -0700, Cesar Philippidis <cesar@codesourcery.com> wrote:
> I've pushed this patch to openacc-gcc-7-branch which teaches the
> gimplifier how to create 0-length arrays mappings for certain pointer
> and reference typed variables. This is necessary to satisfy the implicit
> expectation in OpenACC where the user considers pointer variables used
> like arrays as having array types.
> 
> In OpenACC, pointers variables officially classified as scalar values,
> which requires the content of the scalar to be mapped verbatim onto the
> accelerator (i.e., the host pointer doesn't get remapped to a device
> address on the accelerator). However, as mentioned above, a lot of users
> think that pointers used like arrays are arrays and therefore expect
> things such as
> 
>   int *p = ...
> 
>   #pragma acc enter data copyin (p[0:100])
> 
>   #pragma acc parallel loop
>   for (i = ...)
>   {
>     ... p[i] ...
>   }
> 
> to work as if 'p' were an array.

So you're proposing to change GCC's implementation to deviate from (at
least our current reading of) the OpenACC specification.  I filed
<https://gcc.gnu.org/PR90247> "Reconsider OpenACC implicit data
attributes for pointers".


> Note that this patch does not do
> anything special with dynamic arrays. E.g.
> 
>   int *p = (int *) malloc ...
> 
>   #pragma acc parallel loop
>   for (i = ...)
>   {
>     ... p[i] ...
>   }
> 
> In this case, by virtue of being a 0-length array mapping, p will be
> transferred to the accelerator with the host value of &p[0].

(Shouldn't this be translated into 'NULL' then, according to the OpenMP
"zero-length array section" rules?)

> I've seen
> another compiler which goes through the trouble of mapping the dynamic
> array to the accelerator automatically, but that's beyond the scope of
> this patch.

ACK.


Some quick patch review:

> --- a/gcc/gimplify.c
> +++ b/gcc/gimplify.c
> @@ -6994,7 +6994,12 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags)
>      case ORT_ACC_PARALLEL:
>        rkind = "parallel";

Why would this apply only for the OpenACC 'parallel' construct, but not
for other OpenACC compute constructs?  (This seems inconsistent, just
plain wrong.)

>  
> -      if (is_private)
> +      if (TREE_CODE (type) == REFERENCE_TYPE
> +	  && TREE_CODE (TREE_TYPE (type)) == POINTER_TYPE)
> +	flags |= GOVD_MAP | GOVD_MAP_0LEN_ARRAY;
> +      else if (!lang_GNU_Fortran () && TREE_CODE (type) == POINTER_TYPE)
> +	flags |= GOVD_MAP | GOVD_MAP_0LEN_ARRAY;
> +      else if (is_private)
>  	flags |= GOVD_FIRSTPRIVATE;
>        else if (on_device || declared)
>  	flags |= GOVD_MAP;

The casual reader has no chance to understand that logic, what these two
'if' conditions actually mean.  There are no comments added to the source
code describing what this is about, why these 'if' cases are handled
specially.  There is nothing added to the documentation about the
non-standard behavior (implementation-defined difference from the
official OpenACC specification).

> --- /dev/null
> +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/fp-dyn-arrays.c

This is just one basic execution test case (basically, "my test program X
now works"), which is insufficient for such a change.

> @@ -0,0 +1,42 @@
> +/* Expect dynamic arrays to be mapped on the accelerator via
> +   GOMP_MAP_0LEN_ARRAY.  */
> +
> +#include <stdio.h>
> +#include <stdlib.h>
> +#include <assert.h>
> +
> +int
> +main ()
> +{
> +  const int n = 1000;
> +  int *a, *b, *c, *d, i;
> +
> +  d = (int *) 12345;
> +  a = (int *) malloc (sizeof (int) * n);
> +  b = (int *) malloc (sizeof (int) * n);
> +  c = (int *) malloc (sizeof (int) * n);
> +
> +  for (i = 0; i < n; i++)
> +    {
> +      a[i] = -1;
> +      b[i] = i+1;
> +      c[i] = 2*(i+1);
> +    }
> +
> +#pragma acc enter data create(a[0:n]) copyin(b[:n], c[:n])
> +#pragma acc parallel loop
> +  for (i = 0; i < n; ++i)
> +    {
> +      a[i] = b[i] + c[i] + *((int *)&d);
> +    }
> +#pragma acc exit data copyout(a[0:n]) delete(b[0:n], c[0:n])
> +
> +  for (i = 0; i < n; i++)
> +    assert (a[i] == 3*(i+1) + 12345);
> +
> +  free (a);
> +  free (b);
> +  free (c);
> +
> +  return 0;
> +}


I'm rejecting this, suspending such a change until we get clarification
about the (at least future) intention of the OpenACC specification.

And once we've got that clarified, the patch needs to be re-worked as
indicated.


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/20190425/03950232/attachment.sig>