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