[og7] Enable 0-length array data mappings for implicit data clauses

[Cesar Philippidis]

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. 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]. 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.

Cesar

2017-10-11  Cesar Philippidis  <ce...@codesourcery.com>

	gcc/
	* gimplify.c (oacc_default_clause): Create implicit 0-length
	array data clauses for pointers and reference types.

	libgomp/
	* testsuite/libgomp.oacc-c-c++-common/fp-dyn-arrays.c: New test.


diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 2c10c6433a7..7a9cc241792 100644
--- 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";
 
-      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;
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/fp-dyn-arrays.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/fp-dyn-arrays.c
new file mode 100644
index 00000000000..c57261f2d12
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/fp-dyn-arrays.c
@@ -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;
+}