diff mbox series

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

Message ID 18812b8c-49de-f689-207c-fe97d6397be4@mentor.com
State New
Headers show
Series [OpenACC] Enable 0-length array data mappings for implicit data clauses | expand

Commit Message

Chung-Lin Tang Nov. 6, 2018, 12:49 p.m. UTC 
Hi Thomas, this patch allows the gimplifier to create 0-length array mappings
for certain pointer and reference typed variables. Without this, array usage
of certain pointer variables do not work as usually intended, this is the
original description by Cesar when applied to og7:
https://gcc.gnu.org/ml/gcc-patches/2017-10/msg00673.html

Note that this patch requires this to cleanly apply (still awaiting approval):
https://gcc.gnu.org/ml/gcc-patches/2018-06/msg01911.html

Thomas, since this only touches OpenACC, I suppose you have the powers to approve.

Thanks,
Chung-Lin

2018-11-06  Cesar Philippidis  <cesar@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 mbox series

Patch

diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 9e2b0aa..58ef3de 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -7081,7 +7081,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 0000000..c57261f
--- /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;
+}