diff mbox series

[4/6,og8] Interaction of dynamic/multidimensional arrays with attach/detach.

Message ID 1335fb57035f2e356ecba4f877a82da47d2152fe.1542748807.git.julian@codesourcery.com
State New
Headers show
Series OpenACC attach/detach | expand

Commit Message

Julian Brown Nov. 20, 2018, 9:54 p.m. UTC 
OpenACC multidimensional (or "dynamic") arrays do not seem to fit very
neatly into the attach/detach mechanism described for OpenACC 2.6,
that is if the user tries to use a multidimensional array as a field
in a struct.  This patch disallows that combination, for now at least.
Multidimensional array support in general has been submitted upstream
here but not yet accepted:

https://gcc.gnu.org/ml/gcc-patches/2018-10/msg00937.html

	gcc/
	* omp-low.c (scan_sharing_clauses): Disallow dynamic (multidimensional)
	arrays within structs.

	gcc/testsuite/
	* c-c++-common/goacc/deep-copy-multidim.c: Add test.

	libgomp/
	* target.c (gomp_map_vars_async, gomp_load_image_to_device):
	Zero-initialise do_detach, dynamic_refcount and attach_count in more
	places.
---
 gcc/omp-low.c                                      |   10 +++++-
 .../c-c++-common/goacc/deep-copy-multidim.c        |   32 ++++++++++++++++++++
 libgomp/target.c                                   |    6 ++++
 3 files changed, 47 insertions(+), 1 deletions(-)
 create mode 100644 gcc/testsuite/c-c++-common/goacc/deep-copy-multidim.c
diff mbox series

Patch

diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index e559211..1726451 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -1481,7 +1481,15 @@  scan_sharing_clauses (tree clauses, omp_context *ctx,
 		  t = TREE_TYPE (t);
 		}
 
-	      install_var_field (da_decl, by_ref, 3, ctx);
+	      if (DECL_P (decl))
+		install_var_field (da_decl, by_ref, 3, ctx);
+	      else
+	        {
+		  error_at (OMP_CLAUSE_LOCATION (c),
+			    "dynamic arrays cannot be used within structs");
+		  break;
+		}
+
 	      tree new_var = install_var_local (da_decl, ctx);
 
 	      bool existed = ctx->dynamic_arrays->put (new_var, da_dimensions);
diff --git a/gcc/testsuite/c-c++-common/goacc/deep-copy-multidim.c b/gcc/testsuite/c-c++-common/goacc/deep-copy-multidim.c
new file mode 100644
index 0000000..1696f0c
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/deep-copy-multidim.c
@@ -0,0 +1,32 @@ 
+/* { dg-do compile } */
+
+#include <stdlib.h>
+#include <assert.h>
+
+struct dc
+{
+  int a;
+  int **b;
+};
+
+int
+main ()
+{
+  int n = 100, i, j;
+  struct dc v = { .a = 3 };
+
+  v.b = (int **) malloc (sizeof (int *) * n);
+  for (i = 0; i < n; i++)
+    v.b[i] = (int *) malloc (sizeof (int) * n);
+
+#pragma acc parallel loop copy(v.a, v.b[:n][:n]) /* { dg-error "dynamic arrays cannot be used within structs" } */
+  for (i = 0; i < n; i++)
+    for (j = 0; j < n; j++)
+      v.b[i][j] = v.a + i + j;
+
+  for (i = 0; i < n; i++)
+    for (j = 0; j < n; j++)
+      assert (v.b[i][j] == v.a + i + j);
+
+  return 0;
+}
diff --git a/libgomp/target.c b/libgomp/target.c
index d9d42eb..da51291 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -1484,6 +1484,7 @@  gomp_map_vars_async (struct gomp_device_descr *devicep,
 	     set to false here.  */
 	  tgt->list[i].copy_from = false;
 	  tgt->list[i].always_copy_from = false;
+	  tgt->list[i].do_detach = false;
 
 	  size_t align = (size_t) 1 << (kind >> rshift);
 	  tgt_size = (tgt_size + align - 1) & ~(align - 1);
@@ -1521,6 +1522,8 @@  gomp_map_vars_async (struct gomp_device_descr *devicep,
 
 		  k->tgt = tgt;
 		  k->refcount = 1;
+		  k->dynamic_refcount = 0;
+		  k->attach_count = NULL;
 		  k->link_key = NULL;
 		  tgt_size = (tgt_size + align - 1) & ~(align - 1);
 		  target_row_addr = tgt->tgt_start + tgt_size;
@@ -1532,6 +1535,7 @@  gomp_map_vars_async (struct gomp_device_descr *devicep,
 		    = GOMP_MAP_COPY_FROM_P (kind & typemask);
 		  row_desc->always_copy_from
 		    = GOMP_MAP_ALWAYS_FROM_P (kind & typemask);
+		  row_desc->do_detach = false;
 		  row_desc->offset = 0;
 		  row_desc->length = da->data_row_size;
 
@@ -1839,6 +1843,7 @@  gomp_load_image_to_device (struct gomp_device_descr *devicep, unsigned version,
       k->tgt = tgt;
       k->tgt_offset = target_table[i].start;
       k->refcount = REFCOUNT_INFINITY;
+      k->attach_count = NULL;
       k->link_key = NULL;
       tgt->list[i].key = k;
       tgt->refcount++;
@@ -1873,6 +1878,7 @@  gomp_load_image_to_device (struct gomp_device_descr *devicep, unsigned version,
       k->tgt = tgt;
       k->tgt_offset = target_var->start;
       k->refcount = target_size & link_bit ? REFCOUNT_LINK : REFCOUNT_INFINITY;
+      k->attach_count = NULL;
       k->link_key = NULL;
       tgt->list[i].key = k;
       tgt->refcount++;