[PATCH 5/5] openmp: Support multi-segment noncontiguous descriptors in libgomp
Paul-Antoine Arras <[email protected]> Wed, 5 Aug 2026 22:23:43 +0200
| Newsgroups | gmane.comp.gcc.patches |
|---|---|
| Message-ID | <[email protected]> |
Extend the runtime noncontiguous-array descriptor and its target-side
gather/scatter loops in target.c to walk multiple segments (each
crossing a further pointer indirection) instead of assuming a single
flat set of dimensions, driven by the new nsegments/seg_ndims/seg_nptrs
fields.
Add runtime tests exercising array-shaping casts and array-of-pointers
sections with one, two, and three segments in C and C++.
libgomp/ChangeLog:
* libgomp.h (omp_noncontig_array_desc): Add nsegments, seg_ndims and
seg_nptrs fields; document all fields.
* target.c (omp_target_update_segment): New, extracted and
generalized to walk a single segment of a noncontiguous-array
descriptor.
(gomp_update): Call it once per segment instead of assuming a single
flat set of dimensions.
* testsuite/libgomp.c++/array-section-1.C: New test.
* testsuite/libgomp.c++/array-section-2.C: New test.
* testsuite/libgomp.c-c++-common/array-section-1.c: New test.
* testsuite/libgomp.c-c++-common/array-section-10.c: New test.
* testsuite/libgomp.c-c++-common/array-section-11.c: New test.
* testsuite/libgomp.c-c++-common/array-section-2.c: New test.
* testsuite/libgomp.c-c++-common/array-section-3.c: New test.
* testsuite/libgomp.c-c++-common/array-section-4.c: New test.
* testsuite/libgomp.c-c++-common/array-section-5.c: New test.
* testsuite/libgomp.c-c++-common/array-section-6.c: New test.
* testsuite/libgomp.c-c++-common/array-section-7.c: New test.
* testsuite/libgomp.c-c++-common/array-section-8.c: New test.
* testsuite/libgomp.c-c++-common/array-section-9.c: New test.
---
libgomp/libgomp.h | 24 ++-
libgomp/target.c | 176 ++++++++++++---
.../testsuite/libgomp.c++/array-section-1.C | 204 ++++++++++++++++++
.../testsuite/libgomp.c++/array-section-2.C | 130 +++++++++++
.../libgomp.c-c++-common/array-section-1.c | 181 ++++++++++++++++
.../libgomp.c-c++-common/array-section-10.c | 69 ++++++
.../libgomp.c-c++-common/array-section-11.c | 73 +++++++
.../libgomp.c-c++-common/array-section-2.c | 75 +++++++
.../libgomp.c-c++-common/array-section-3.c | 135 ++++++++++++
.../libgomp.c-c++-common/array-section-4.c | 134 ++++++++++++
.../libgomp.c-c++-common/array-section-5.c | 136 ++++++++++++
.../libgomp.c-c++-common/array-section-6.c | 75 +++++++
.../libgomp.c-c++-common/array-section-7.c | 104 +++++++++
.../libgomp.c-c++-common/array-section-8.c | 87 ++++++++
.../libgomp.c-c++-common/array-section-9.c | 120 +++++++++++
15 files changed, 1686 insertions(+), 37 deletions(-)
create mode 100644 libgomp/testsuite/libgomp.c++/array-section-1.C
create mode 100644 libgomp/testsuite/libgomp.c++/array-section-2.C
create mode 100644 libgomp/testsuite/libgomp.c-c++-common/array-section-1.c
create mode 100644 libgomp/testsuite/libgomp.c-c++-common/array-section-10.c
create mode 100644 libgomp/testsuite/libgomp.c-c++-common/array-section-11.c
create mode 100644 libgomp/testsuite/libgomp.c-c++-common/array-section-2.c
create mode 100644 libgomp/testsuite/libgomp.c-c++-common/array-section-3.c
create mode 100644 libgomp/testsuite/libgomp.c-c++-common/array-section-4.c
create mode 100644 libgomp/testsuite/libgomp.c-c++-common/array-section-5.c
create mode 100644 libgomp/testsuite/libgomp.c-c++-common/array-section-6.c
create mode 100644 libgomp/testsuite/libgomp.c-c++-common/array-section-7.c
create mode 100644 libgomp/testsuite/libgomp.c-c++-common/array-section-8.c
create mode 100644 libgomp/testsuite/libgomp.c-c++-common/array-section-9.c
diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h
index 9f13f9e32c2..eadecbc0055 100644
--- a/libgomp/libgomp.h
+++ b/libgomp/libgomp.h
@@ -1385,13 +1385,23 @@ struct target_mem_desc {
omp-low.cc:omp_noncontig_descriptor_type. */
typedef struct {
- size_t ndims;
- size_t elemsize;
- size_t span;
- size_t *dim;
- size_t *index;
- size_t *length;
- size_t *stride;
+ size_t ndims; /* Number of array dimensions */
+ size_t elemsize; /* Byte size of each innermost element */
+ size_t
+ span; /* Byte distance between two logically-adjacent elements. Defaults to
+ elemsize but can be more if stride is specified. */
+ size_t *dim; /* Per-dimension total extent */
+ size_t *index; /* Per-dimension start offset */
+ size_t *length; /* Per-dimension count of selected elements (1 for a fixed
+ index, more for a range) */
+ size_t *stride; /* Per-dimension element stride (defaults to 1) */
+ size_t nsegments; /* Number of sets of contiguous dimensions
+ (= 1 + number of indirections) */
+ size_t *seg_ndims; /* Per-segment dimension count */
+ size_t *seg_nptrs; /* For each segment junction, number of pointers
+ selected by the indirection between segment k and k+1
+ (1 for a fixed index, N for a strided/ranged dimension
+ in segment k) */
} omp_noncontig_array_desc;
diff --git a/libgomp/target.c b/libgomp/target.c
index 859e2de1847..6c5650d7f42 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -2693,6 +2693,103 @@ omp_target_memcpy_rect_worker (void *, const void *, size_t, size_t, int,
struct gomp_device_descr *, size_t *tmp_size,
void **tmp);
+/* Transfer BASE to or from the device for one segment of DESC's pointer-
+ indirected dimensions, recursing through each indirection until the
+ innermost (contiguous) segment, where the actual rectangular memcpy
+ happens. SEG_START holds each segment's starting index into DESC's flat
+ per-dimension arrays (dim/index/length/stride). SEG is the segment
+ currently being processed. BASE is the host address this segment
+ indexes into -- the original host pointer for segment 0, or the pointee
+ reached through the previous segment's selected pointer otherwise. */
+
+static void
+omp_target_update_segment (omp_noncontig_array_desc *desc,
+ struct gomp_device_descr *devicep, bool to_not_from,
+ size_t seg_start[], int seg, void *base)
+{
+ size_t start = seg_start[seg];
+ size_t ndims = desc->seg_ndims[seg];
+
+ /* Inner segment: do the actual memory transfer. */
+ if (seg == desc->nsegments - 1)
+ {
+ /* Get device address. */
+ struct splay_tree_key_s cur_node;
+ cur_node.host_start = (uintptr_t) base;
+ cur_node.host_end = cur_node.host_start + 1;
+ splay_tree_key tn = splay_tree_lookup (&devicep->mem_map, &cur_node);
+ if (tn == NULL)
+ {
+ gomp_mutex_unlock (&devicep->lock);
+ gomp_error ("pointer not mapped");
+ return;
+ }
+
+ /* Do rectangular memcpy. */
+ void *devaddr = (void *) (tn->tgt->tgt_start + tn->tgt_offset);
+ size_t inner_seg = desc->nsegments - 1;
+ size_t inner_dim = desc->ndims - desc->seg_ndims[inner_seg];
+ size_t tmp_size = 0;
+ void *tmp = NULL;
+ if (to_not_from)
+ omp_target_memcpy_rect_worker (
+ devaddr, base, desc->elemsize, desc->span,
+ desc->seg_ndims[desc->nsegments - 1], &desc->length[inner_dim],
+ &desc->stride[inner_dim], &desc->index[inner_dim],
+ &desc->index[inner_dim], &desc->dim[inner_dim], &desc->dim[inner_dim],
+ devicep, NULL, &tmp_size, &tmp);
+ else
+ omp_target_memcpy_rect_worker (
+ base, devaddr, desc->elemsize, desc->span,
+ desc->seg_ndims[desc->nsegments - 1], &desc->length[inner_dim],
+ &desc->stride[inner_dim], &desc->index[inner_dim],
+ &desc->index[inner_dim], &desc->dim[inner_dim], &desc->dim[inner_dim],
+ NULL, devicep, &tmp_size, &tmp);
+ return;
+ }
+
+ /* Intermediate segment: compute position of each pointer and thread pointee
+ through recursion. */
+ for (size_t j = 0; j < desc->seg_nptrs[seg]; j++)
+ {
+ size_t rem = j;
+ size_t idx[ndims];
+ for (size_t dd = 0; dd < ndims; dd++)
+ {
+ size_t d = start + ndims - 1 - dd;
+ size_t kk = rem % desc->length[d];
+ rem /= desc->length[d];
+ idx[d - start] = desc->index[d] + kk * desc->stride[d];
+ }
+ size_t pos = 0;
+ for (size_t d = start; d < start + ndims; d++)
+ pos = pos * desc->dim[d] + idx[d - start];
+
+ void **host_ptr = (void **) base + pos;
+ void *inner_base = *host_ptr;
+
+ omp_target_update_segment (desc, devicep, to_not_from, seg_start, seg + 1,
+ inner_base);
+ }
+}
+
+/* Perform a "target update" on DEVICEP, copying each of the MAPNUM regions
+ described by HOSTADDRS/SIZES/KINDS (KINDS encoded per SHORT_MAPKIND)
+ between host and device, using whichever device mapping is already
+ recorded in DEVICEP's memory map. A region not currently mapped is
+ silently skipped, unless its map kind requires the region to be present,
+ in which case that is a fatal error.
+
+ A GOMP_MAP_TO_GRID/GOMP_MAP_FROM_GRID entry describes a non-contiguous
+ (reshaped or strided) array section rather than a single flat region;
+ HOSTADDRS[I + 1] then points to the omp_noncontig_array_desc describing
+ its dimensions, which is copied one dimension at a time via
+ omp_target_memcpy_rect_worker and consumes the following entry as well.
+
+ For an ordinary region whose underlying object currently has pointers
+ being attached, only the individual pointer-sized slots not in the
+ middle of being attached are updated. */
+
static void
gomp_update (struct gomp_device_descr *devicep, size_t mapnum, void **hostaddrs,
size_t *sizes, void *kinds, bool short_mapkind)
@@ -2729,39 +2826,58 @@ gomp_update (struct gomp_device_descr *devicep, size_t mapnum, void **hostaddrs,
omp_noncontig_array_desc *desc
= (omp_noncontig_array_desc *) hostaddrs[i + 1];
size_t bias = sizes[i + 1];
- cur_node.host_start = (uintptr_t) hostaddrs[i] + bias;
- cur_node.host_end = cur_node.host_start + sizes[i];
- splay_tree_key n = splay_tree_lookup (&devicep->mem_map, &cur_node);
- if (n)
+
+ if (desc->nsegments == 1)
{
- if (n->aux && n->aux->attach_count)
+ cur_node.host_start = (uintptr_t) hostaddrs[i] + bias;
+ cur_node.host_end = cur_node.host_start + sizes[i];
+ splay_tree_key n
+ = splay_tree_lookup (&devicep->mem_map, &cur_node);
+ if (n)
{
- gomp_mutex_unlock (&devicep->lock);
- gomp_error ("noncontiguous update with attached pointers");
- return;
+ if (n->aux && n->aux->attach_count)
+ {
+ gomp_mutex_unlock (&devicep->lock);
+ gomp_error (
+ "noncontiguous update with attached pointers");
+ return;
+ }
+ void *devaddr
+ = (void *) (n->tgt->tgt_start + n->tgt_offset
+ + cur_node.host_start - n->host_start - bias);
+ size_t tmp_size = 0;
+ void *tmp = NULL;
+ if ((kind & typemask) == GOMP_MAP_TO_GRID)
+ omp_target_memcpy_rect_worker (devaddr, hostaddrs[i],
+ desc->elemsize, desc->span,
+ desc->ndims, desc->length,
+ desc->stride, desc->index,
+ desc->index, desc->dim,
+ desc->dim, devicep, NULL,
+ &tmp_size, &tmp);
+ else
+ omp_target_memcpy_rect_worker (hostaddrs[i], devaddr,
+ desc->elemsize, desc->span,
+ desc->ndims, desc->length,
+ desc->stride, desc->index,
+ desc->index, desc->dim,
+ desc->dim, NULL, devicep,
+ &tmp_size, &tmp);
}
- void *devaddr = (void *) (n->tgt->tgt_start + n->tgt_offset
- + cur_node.host_start
- - n->host_start
- - bias);
- size_t tmp_size = 0;
- void *tmp = NULL;
- if ((kind & typemask) == GOMP_MAP_TO_GRID)
- omp_target_memcpy_rect_worker (devaddr, hostaddrs[i],
- desc->elemsize, desc->span,
- desc->ndims, desc->length,
- desc->stride, desc->index,
- desc->index, desc->dim,
- desc->dim, devicep,
- NULL, &tmp_size, &tmp);
- else
- omp_target_memcpy_rect_worker (hostaddrs[i], devaddr,
- desc->elemsize, desc->span,
- desc->ndims, desc->length,
- desc->stride, desc->index,
- desc->index, desc->dim,
- desc->dim, NULL,
- devicep, &tmp_size, &tmp);
+ }
+ else
+ {
+ assert (desc->nsegments != 0);
+
+ /* Compute the start dimension for each segment. */
+ size_t seg_start[desc->nsegments];
+ seg_start[0] = 0;
+ for (int s = 1; s < desc->nsegments; s++)
+ seg_start[s] = seg_start[s - 1] + desc->seg_ndims[s - 1];
+
+ omp_target_update_segment (desc, devicep,
+ (kind & typemask) == GOMP_MAP_TO_GRID,
+ seg_start, 0, hostaddrs[i]);
}
i++;
}
diff --git a/libgomp/testsuite/libgomp.c++/array-section-1.C b/libgomp/testsuite/libgomp.c++/array-section-1.C
new file mode 100644
index 00000000000..cdd396efbd8
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/array-section-1.C
@@ -0,0 +1,204 @@
+// { dg-do run { target offload_device_nonshared_as } }
+
+/* Noncontiguous "target update" through an array of pointers accessed via
+ a C++ reference + array-shaped variant. */
+
+#include <string.h>
+#include <assert.h>
+#include <stdlib.h>
+
+#define DIM1 5
+#define DIM2 10
+#define ROWS 6
+#define COLS 6
+
+void
+test_fixed_index ()
+{
+ int *x[DIM1];
+ int *(&x_ref)[DIM1] = x;
+ int i, j;
+
+ for (i = 0; i < DIM1; i++)
+ x[i] = (int *) calloc (DIM2, sizeof (int));
+
+#pragma omp target enter data map(alloc: x_ref[0:DIM1])
+#pragma omp target enter data map(to: x_ref[2][:DIM2])
+
+ for (j = 0; j < DIM2; j++)
+ x_ref[2][j] = j + 1;
+
+#pragma omp target update to(x_ref[2][:DIM2])
+
+#pragma omp target map(alloc: x_ref)
+ {
+ for (int j = 0; j < DIM2; j++)
+ x_ref[2][j] += 1000;
+ }
+
+ memset (x[2], 0, DIM2 * sizeof (int));
+
+#pragma omp target exit data map(from: x_ref[2][:DIM2])
+#pragma omp target exit data map(release: x_ref[0:DIM1])
+
+ for (j = 0; j < DIM2; j++)
+ assert (x[2][j] == j + 1 + 1000);
+
+ for (i = 0; i < DIM1; i++)
+ free (x[i]);
+}
+
+void
+test_strided_range ()
+{
+ int *x[DIM1];
+ int *(&x_ref)[DIM1] = x;
+ int i, j;
+
+ for (i = 0; i < DIM1; i++)
+ x[i] = (int *) calloc (DIM2, sizeof (int));
+
+#pragma omp target enter data map(alloc: x_ref[0:DIM1])
+
+ for (i = 0; i < DIM1; i++)
+ {
+#pragma omp target enter data map(to: x_ref[i][:DIM2])
+ }
+
+ for (i = 0; i < DIM1; i++)
+ for (j = 0; j < DIM2; j++)
+ if (i % 2 == 0)
+ x_ref[i][j] = 10 * i + j;
+
+#pragma omp target update to(x_ref[0:3:2][:DIM2])
+
+#pragma omp target map(alloc: x_ref)
+ {
+ for (int i = 0; i < DIM1; i += 2)
+ for (int j = 0; j < DIM2; j++)
+ x_ref[i][j] += 1000;
+ }
+
+ for (i = 0; i < DIM1; i++)
+ memset (x[i], 0, DIM2 * sizeof (int));
+
+ for (i = 0; i < DIM1; i++)
+ {
+#pragma omp target exit data map(from: x_ref[i][:DIM2])
+ }
+
+#pragma omp target exit data map(release: x_ref[0:DIM1])
+
+ for (i = 0; i < DIM1; i++)
+ for (j = 0; j < DIM2; j++)
+ if (i % 2 == 0)
+ assert (x[i][j] == 10 * i + j + 1000);
+ else
+ assert (x[i][j] == 0);
+
+ for (i = 0; i < DIM1; i++)
+ free (x[i]);
+}
+
+void
+test_update_from ()
+{
+ int *x[DIM1];
+ int *(&x_ref)[DIM1] = x;
+ int i, j;
+
+ for (i = 0; i < DIM1; i++)
+ x[i] = (int *) calloc (DIM2, sizeof (int));
+
+#pragma omp target enter data map(alloc: x_ref[0:DIM1])
+#pragma omp target enter data map(to: x_ref[3][:DIM2])
+
+ for (j = 0; j < DIM2; j++)
+ x_ref[3][j] = 100 + j;
+
+#pragma omp target update to(x_ref[3][:DIM2])
+
+#pragma omp target map(alloc: x_ref)
+ {
+ for (int j = 0; j < DIM2; j++)
+ x_ref[3][j] += 1000;
+ }
+
+ for (j = 0; j < DIM2; j++)
+ x_ref[3][j] = -1;
+
+#pragma omp target update from(x_ref[3][:DIM2])
+
+ for (j = 0; j < DIM2; j++)
+ assert (x[3][j] == 100 + j + 1000);
+
+#pragma omp target exit data map(delete: x_ref[3][:DIM2])
+#pragma omp target exit data map(release: x_ref[0:DIM1])
+
+ for (i = 0; i < DIM1; i++)
+ free (x[i]);
+}
+
+void
+test_shape_cast ()
+{
+ int *x[DIM1];
+ int *(&x_ref)[DIM1] = x;
+ int i, r, c;
+
+ for (i = 0; i < DIM1; i++)
+ x[i] = (int *) calloc (ROWS * COLS, sizeof (int));
+
+#pragma omp target enter data map(alloc: x_ref[0:DIM1])
+
+ for (i = 0; i < DIM1; i++)
+ {
+#pragma omp target enter data map(to: x_ref[i][:ROWS*COLS])
+ }
+
+ for (r = 0; r < ROWS; r++)
+ for (c = 0; c < COLS; c++)
+ x_ref[1][r * COLS + c] = 100 * r + c;
+
+#pragma omp target update to((([ROWS][COLS]) x_ref[1])[1:3:2][0:2])
+
+#pragma omp target map(alloc: x_ref)
+ {
+ for (int r = 1; r <= 5; r += 2)
+ for (int c = 0; c < 2; c++)
+ x_ref[1][r * COLS + c] += 1000;
+ }
+
+ for (i = 0; i < DIM1; i++)
+ memset (x[i], 0, ROWS * COLS * sizeof (int));
+
+ for (i = 0; i < DIM1; i++)
+ {
+#pragma omp target exit data map(from: x_ref[i][:ROWS*COLS])
+ }
+
+#pragma omp target exit data map(release: x_ref[0:DIM1])
+
+ for (i = 0; i < DIM1; i++)
+ for (r = 0; r < ROWS; r++)
+ for (c = 0; c < COLS; c++)
+ {
+ int expected = 0;
+ if (i == 1 && (r == 1 || r == 3 || r == 5) && c < 2)
+ expected = 100 * r + c + 1000;
+ assert (x[i][r * COLS + c] == expected);
+ }
+
+ for (i = 0; i < DIM1; i++)
+ free (x[i]);
+}
+
+int
+main ()
+{
+ test_fixed_index ();
+ test_strided_range ();
+ test_update_from ();
+ test_shape_cast ();
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c++/array-section-2.C b/libgomp/testsuite/libgomp.c++/array-section-2.C
new file mode 100644
index 00000000000..c37f3200cce
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/array-section-2.C
@@ -0,0 +1,130 @@
+// { dg-do run { target offload_device_nonshared_as } }
+
+/* Test noncontiguous "target update" with three segments (two pointer indirections)
+ and two dimensions per segment, along with C++ templates and references. */
+
+#include <assert.h>
+#include <stdlib.h>
+
+#define DIM1A 2
+#define DIM1B 3
+#define ROWS1 3
+#define NCOLS 3
+#define ROWS2 3
+#define ROWLEN 20
+
+template<typename T>
+static T
+encode (int i, int j, int r1, int c1, int r2, int k)
+{
+ return (T) (i * 1000000 + j * 100000 + r1 * 10000 + c1 * 1000
+ + r2 * 100 + k);
+}
+
+template<typename T>
+void
+test ()
+{
+ typedef T (*p2_t)[ROWLEN];
+ typedef p2_t (*p1_t)[NCOLS];
+
+ p1_t arr[DIM1A][DIM1B];
+ p1_t (&arr_ref)[DIM1A][DIM1B] = arr;
+ int i, j, r1, c1, r2, k;
+
+ for (i = 0; i < DIM1A; i++)
+ for (j = 0; j < DIM1B; j++)
+ {
+ arr_ref[i][j] = (p1_t) malloc (ROWS1 * NCOLS * sizeof (p2_t));
+ for (r1 = 0; r1 < ROWS1; r1++)
+ for (c1 = 0; c1 < NCOLS; c1++)
+ arr_ref[i][j][r1][c1]
+ = (p2_t) malloc (ROWS2 * ROWLEN * sizeof (T));
+ }
+
+ for (i = 0; i < DIM1A; i++)
+ for (j = 0; j < DIM1B; j++)
+ for (r1 = 0; r1 < ROWS1; r1++)
+ for (c1 = 0; c1 < NCOLS; c1++)
+ for (r2 = 0; r2 < ROWS2; r2++)
+ for (k = 0; k < ROWLEN; k++)
+ arr_ref[i][j][r1][c1][r2][k] = encode<T> (i, j, r1, c1, r2, k);
+
+ for (i = 0; i < DIM1A; i++)
+ for (j = 0; j < DIM1B; j++)
+ for (r1 = 0; r1 < ROWS1; r1++)
+ for (c1 = 0; c1 < NCOLS; c1++)
+ {
+#pragma omp target enter data map(to: arr_ref[i][j][r1][c1][:ROWS2])
+ }
+
+ /* Corrupt the whole of ARR_REF[1]'s data, so that precisely matching
+ the six-dimensional touched set below after the update proves each
+ dimension of each segment was walked correctly. */
+ for (j = 0; j < DIM1B; j++)
+ for (r1 = 0; r1 < ROWS1; r1++)
+ for (c1 = 0; c1 < NCOLS; c1++)
+ for (r2 = 0; r2 < ROWS2; r2++)
+ for (k = 0; k < ROWLEN; k++)
+ arr_ref[1][j][r1][c1][r2][k] = (T) -1;
+
+#pragma omp target update from(arr_ref[1][0:2:2][0:2:2][0:2:2][0:2][3:5:2])
+
+ for (i = 0; i < DIM1A; i++)
+ for (j = 0; j < DIM1B; j++)
+ for (r1 = 0; r1 < ROWS1; r1++)
+ for (c1 = 0; c1 < NCOLS; c1++)
+ for (r2 = 0; r2 < ROWS2; r2++)
+ for (k = 0; k < ROWLEN; k++)
+ {
+ T expected;
+ if (i == 1)
+ {
+ int j_touched = (j == 0 || j == 2);
+ int r1_touched = (r1 == 0 || r1 == 2);
+ int c1_touched = (c1 == 0 || c1 == 2);
+ int r2_touched = (r2 == 0 || r2 == 1);
+ int k_touched
+ = (k == 3 || k == 5 || k == 7 || k == 9 || k == 11);
+ if (j_touched && r1_touched && c1_touched
+ && r2_touched && k_touched)
+ /* Restored from the device -- proves this exact
+ combination of dimensions across all three
+ segments was selected. */
+ expected = encode<T> (i, j, r1, c1, r2, k);
+ else
+ /* Still corrupted -- proves the update did not
+ touch anything outside the requested set. */
+ expected = (T) -1;
+ }
+ else
+ /* Untouched by either the corruption or the update. */
+ expected = encode<T> (i, j, r1, c1, r2, k);
+ assert (arr_ref[i][j][r1][c1][r2][k] == expected);
+ }
+
+ for (i = 0; i < DIM1A; i++)
+ for (j = 0; j < DIM1B; j++)
+ for (r1 = 0; r1 < ROWS1; r1++)
+ for (c1 = 0; c1 < NCOLS; c1++)
+ {
+#pragma omp target exit data map(delete: arr_ref[i][j][r1][c1][:ROWS2])
+ }
+
+ for (i = 0; i < DIM1A; i++)
+ for (j = 0; j < DIM1B; j++)
+ {
+ for (r1 = 0; r1 < ROWS1; r1++)
+ for (c1 = 0; c1 < NCOLS; c1++)
+ free (arr_ref[i][j][r1][c1]);
+ free (arr_ref[i][j]);
+ }
+}
+
+int
+main ()
+{
+ test<int> ();
+ test<long> ();
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/array-section-1.c b/libgomp/testsuite/libgomp.c-c++-common/array-section-1.c
new file mode 100644
index 00000000000..514367f16a5
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/array-section-1.c
@@ -0,0 +1,181 @@
+/* { dg-do run { target offload_device_nonshared_as } } */
+
+/* Noncontiguous "target update" through one level of pointer indirection,
+ i.e. an array of dynamically-allocated pointers. */
+
+#include <string.h>
+#include <assert.h>
+#include <stdlib.h>
+
+#define DIM1 5
+#define DIM2 10
+
+int main()
+{
+ int *x[DIM1];
+ int i, j;
+
+ for (i = 0; i < DIM1; i++)
+ x[i] = (int *)calloc (DIM2, sizeof (int));
+
+#pragma omp target enter data map(alloc: x[0:DIM1])
+
+ /* Case 1: Fixed-index update to device through one pointer. */
+
+ for (i = 0; i < DIM1; i++)
+ {
+#pragma omp target enter data map(to: x[i][:DIM2])
+ }
+
+ for (j = 0; j < DIM2; j++)
+ x[2][j] = j + 1;
+
+#pragma omp target update to(x[2][:DIM2])
+
+#pragma omp target map(alloc: x)
+ {
+ for (int j = 0; j < DIM2; j++)
+ x[2][j] += 1000;
+ }
+
+ /* Erase the host copies before pulling back from the device, so the
+ assertions below only see what actually made it across. */
+ for (i = 0; i < DIM1; i++)
+ memset (x[i], 0, DIM2 * sizeof (int));
+
+ for (i = 0; i < DIM1; i++)
+ {
+#pragma omp target exit data map(from: x[i][:DIM2])
+ }
+
+ for (i = 0; i < DIM1; i++)
+ for (j = 0; j < DIM2; j++)
+ if (i == 2)
+ assert (x[i][j] == j + 1 + 1000);
+ else
+ assert (x[i][j] == 0);
+
+ /* Case 2: Range of pointers: a single directive walks through several
+ selected pointers, each contributing its own contiguous run. */
+
+ for (i = 0; i < DIM1; i++)
+ memset (x[i], 0, DIM2 * sizeof (int));
+
+ for (i = 0; i < DIM1; i++)
+ {
+#pragma omp target enter data map(to: x[i][:DIM2])
+ }
+
+ for (i = 1; i <= 3; i++)
+ for (j = 0; j < DIM2; j++)
+ x[i][j] = 10 * i + j;
+
+#pragma omp target update to(x[1:3][:DIM2])
+
+#pragma omp target map(alloc: x)
+ {
+ for (int i = 1; i <= 3; i++)
+ for (int j = 0; j < DIM2; j++)
+ x[i][j] += 1000;
+ }
+
+ for (i = 0; i < DIM1; i++)
+ memset (x[i], 0, DIM2 * sizeof (int));
+
+ for (i = 0; i < DIM1; i++)
+ {
+#pragma omp target exit data map(from: x[i][:DIM2])
+ }
+
+ for (i = 0; i < DIM1; i++)
+ for (j = 0; j < DIM2; j++)
+ if (i >= 1 && i <= 3)
+ assert (x[i][j] == 10 * i + j + 1000);
+ else
+ assert (x[i][j] == 0);
+
+ /* Case 3: Strided range of pointers. */
+
+ for (i = 0; i < DIM1; i++)
+ memset (x[i], 0, DIM2 * sizeof (int));
+
+ for (i = 0; i < DIM1; i++)
+ {
+#pragma omp target enter data map(to: x[i][:DIM2])
+ }
+
+ for (i = 0; i < DIM1; i++)
+ for (j = 0; j < DIM2; j++)
+ if (i % 2 == 0)
+ x[i][j] = 10 * i + j;
+
+#pragma omp target update to(x[0:3:2][:DIM2])
+
+#pragma omp target map(alloc: x)
+ {
+ for (int i = 0; i < DIM1; i += 2)
+ for (int j = 0; j < DIM2; j++)
+ x[i][j] += 1000;
+ }
+
+ for (i = 0; i < DIM1; i++)
+ memset (x[i], 0, DIM2 * sizeof (int));
+
+ for (i = 0; i < DIM1; i++)
+ {
+#pragma omp target exit data map(from: x[i][:DIM2])
+ }
+
+ for (i = 0; i < DIM1; i++)
+ for (j = 0; j < DIM2; j++)
+ if (i%2 == 0)
+ assert (x[i][j] == 10 * i + j + 1000);
+ else
+ assert (x[i][j] == 0);
+
+ /* Case 4: Update from device through the indirection + target region
+ computation. */
+
+ for (i = 0; i < DIM1; i++)
+ memset (x[i], 0, DIM2 * sizeof (int));
+
+ for (i = 0; i < DIM1; i++)
+ {
+#pragma omp target enter data map(to: x[i][:DIM2])
+ }
+
+ for (j = 0; j < DIM2; j++)
+ x[3][j] = 100 + j;
+
+#pragma omp target update to(x[3][:DIM2])
+
+#pragma omp target map(alloc: x)
+ {
+ for (int j = 0; j < DIM2; j++)
+ x[3][j] += 1000;
+ }
+
+ for (j = 0; j < DIM2; j++)
+ x[3][j] = -1;
+
+#pragma omp target update from(x[3][:DIM2])
+
+ for (i = 0; i < DIM1; i++)
+ for (j = 0; j < DIM2; j++)
+ if (i == 3)
+ assert (x[i][j] == 100 + j + 1000);
+ else
+ assert (x[i][j] == 0);
+
+ for (i = 0; i < DIM1; i++)
+ {
+#pragma omp target exit data map(delete: x[i][:DIM2])
+ }
+
+#pragma omp target exit data map(release: x[0:DIM1])
+
+ for (i = 0; i < DIM1; i++)
+ free (x[i]);
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/array-section-10.c b/libgomp/testsuite/libgomp.c-c++-common/array-section-10.c
new file mode 100644
index 00000000000..531f3a641ca
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/array-section-10.c
@@ -0,0 +1,69 @@
+/* { dg-do run { target offload_device_nonshared_as } } */
+
+/* Check that data transferred through noncontiguous target update can be
+ accessed from the device. */
+
+#include <assert.h>
+#include <stdlib.h>
+
+#define DIM1 5
+#define DIM2 10
+
+int
+main ()
+{
+ int *x[DIM1];
+ int i, j;
+
+ for (i = 0; i < DIM1; i++)
+ x[i] = (int *) calloc (DIM2, sizeof (int));
+
+ for (i = 0; i < DIM1; i++)
+ for (j = 0; j < DIM2; j++)
+ x[i][j] = 100 * i + j;
+
+#pragma omp target enter data map(alloc: x[0:DIM1])
+
+ for (i = 0; i < DIM1; i++)
+ {
+#pragma omp target enter data map(to: x[i][:DIM2])
+ }
+
+ for (j = 0; j < DIM2; j += 2)
+ x[1][j] = 1000 + j;
+
+#pragma omp target update to(x[1][0:DIM2/2:2])
+
+ /* Corrupt the host copy of the whole row, so the assertions at the
+ end can only pass if the values genuinely came from the device --
+ both from the strided update above and from the target region
+ below. */
+ for (j = 0; j < DIM2; j++)
+ x[1][j] = -1;
+
+#pragma omp target map(alloc: x)
+ {
+ for (int j = 0; j < DIM2; j++)
+ x[1][j] += 1;
+ }
+
+#pragma omp target update from(x[1][:DIM2])
+
+ for (j = 0; j < DIM2; j++)
+ {
+ int expected = (j % 2 == 0) ? (1000 + j + 1) : (100 * 1 + j + 1);
+ assert (x[1][j] == expected);
+ }
+
+ for (i = 0; i < DIM1; i++)
+ {
+#pragma omp target exit data map(delete: x[i][:DIM2])
+ }
+
+#pragma omp target exit data map(delete: x[0:DIM1])
+
+ for (i = 0; i < DIM1; i++)
+ free (x[i]);
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/array-section-11.c b/libgomp/testsuite/libgomp.c-c++-common/array-section-11.c
new file mode 100644
index 00000000000..82ed09d9314
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/array-section-11.c
@@ -0,0 +1,73 @@
+/* { dg-do run { target offload_device_nonshared_as } } */
+/* { dg-xfail-run-if "" { offload_device_nonshared_as } } */
+
+/* Two levels of pointer indirection dereferenced inside a target region cause
+ an illegal memory access, even without noncontiguous updates. */
+
+#include <assert.h>
+#include <stdlib.h>
+
+#define DIM1 4
+#define DIM2 3
+#define N 20
+
+int
+main ()
+{
+ int **arr[DIM1];
+ int i, j, k;
+
+ for (i = 0; i < DIM1; i++)
+ {
+ arr[i] = (int **) malloc (DIM2 * sizeof (int *));
+ for (j = 0; j < DIM2; j++)
+ arr[i][j] = (int *) calloc (N, sizeof (int));
+ }
+
+ /* Manual deep copy. */
+#pragma omp target enter data map(alloc: arr[0:DIM1])
+ for (i = 0; i < DIM1; i++)
+ {
+#pragma omp target enter data map(alloc: arr[i][0:DIM2])
+ }
+ for (i = 0; i < DIM1; i++)
+ for (j = 0; j < DIM2; j++)
+ {
+#pragma omp target enter data map(to: arr[i][j][0:N])
+ }
+
+ /* Crash here: dereferencing two levels of indirection to write
+ data from target code. */
+#pragma omp target map(alloc: arr)
+ {
+ for (int i = 0; i < DIM1; i++)
+ for (int j = 0; j < DIM2; j++)
+ for (int k = 0; k < N; k++)
+ arr[i][j][k] = i * 10000 + j * 100 + k;
+ }
+
+ for (i = 0; i < DIM1; i++)
+ for (j = 0; j < DIM2; j++)
+ {
+#pragma omp target exit data map(release: arr[i][j]) map(from: arr[i][j][0:N])
+ }
+ for (i = 0; i < DIM1; i++)
+ {
+#pragma omp target exit data map(release: arr[i], arr[i][0:DIM2])
+ }
+#pragma omp target exit data map(release: arr, arr[0:DIM1])
+
+ for (i = 0; i < DIM1; i++)
+ for (j = 0; j < DIM2; j++)
+ for (k = 0; k < N; k++)
+ assert (arr[i][j][k] == i * 10000 + j * 100 + k);
+
+ for (i = 0; i < DIM1; i++)
+ {
+ for (j = 0; j < DIM2; j++)
+ free (arr[i][j]);
+ free (arr[i]);
+ }
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/array-section-2.c b/libgomp/testsuite/libgomp.c-c++-common/array-section-2.c
new file mode 100644
index 00000000000..a8982133a95
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/array-section-2.c
@@ -0,0 +1,75 @@
+/* { dg-do run { target offload_device_nonshared_as } } */
+
+/* Noncontiguous "target update" through a 2D array of pointers where the
+ pointers themselves are the mapped payload rather than a point of indirection
+ -- one segment, two dimensions, with the final pointer left un-dereferenced.
+ */
+
+#include <assert.h>
+
+#define DIM1 5
+#define DIM2 5
+
+int main()
+{
+ int *y[DIM1][DIM2];
+ int markers[2];
+ int i, j;
+
+ /* Case 1: update to device, a partial range of one fixed row. */
+
+ for (i = 0; i < DIM1; i++)
+ for (j = 0; j < DIM2; j++)
+ y[i][j] = 0;
+
+#pragma omp target enter data map(to: y)
+
+ for (j = 0; j < 2; j++)
+ y[0][j] = &markers[j];
+
+#pragma omp target update to(y[0][:2])
+
+ /* Erase the host copy before pulling back from the device, so the
+ assertions below only see what actually made it across. */
+ for (i = 0; i < DIM1; i++)
+ for (j = 0; j < DIM2; j++)
+ y[i][j] = 0;
+
+#pragma omp target exit data map(from: y)
+
+ for (i = 0; i < DIM1; i++)
+ for (j = 0; j < DIM2; j++)
+ if (i == 0 && j < 2)
+ assert (y[i][j] == &markers[j]);
+ else
+ assert (y[i][j] == 0);
+
+ /* Case 2: update from device, same shape, on a different row. */
+
+ for (i = 0; i < DIM1; i++)
+ for (j = 0; j < DIM2; j++)
+ y[i][j] = 0;
+
+#pragma omp target enter data map(to: y)
+
+ for (j = 0; j < 2; j++)
+ y[3][j] = &markers[j];
+
+#pragma omp target update to(y[3][:2])
+
+ for (j = 0; j < 2; j++)
+ y[3][j] = 0;
+
+#pragma omp target update from(y[3][:2])
+
+ for (i = 0; i < DIM1; i++)
+ for (j = 0; j < DIM2; j++)
+ if (i == 3 && j < 2)
+ assert (y[i][j] == &markers[j]);
+ else
+ assert (y[i][j] == 0);
+
+#pragma omp target exit data map(delete: y)
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/array-section-3.c b/libgomp/testsuite/libgomp.c-c++-common/array-section-3.c
new file mode 100644
index 00000000000..10eb13d600a
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/array-section-3.c
@@ -0,0 +1,135 @@
+/* { dg-do run { target offload_device_nonshared_as } } */
+
+/* Test noncontiguous "target update" through one level of array-shaped pointer
+ indirection, with two dimensions on the array-of-pointers side.
+ Also check the transferred data can be accessed from device code. */
+
+#include <string.h>
+#include <assert.h>
+#include <stdlib.h>
+
+#define DIM1 3
+#define DIM2 4
+#define ROWS 6
+#define COLS 6
+
+int main()
+{
+ int *x[DIM1][DIM2];
+ int i, j, r, c;
+
+ for (i = 0; i < DIM1; i++)
+ for (j = 0; j < DIM2; j++)
+ x[i][j] = (int *) calloc (ROWS * COLS, sizeof (int));
+
+#pragma omp target enter data map(alloc: x[0:DIM1])
+
+ /* Case 1: update to device. Two fixed indices select a single pointer
+ from the 2D array-of-pointers (segment 0); the selected pointer
+ contributes a 2D sub-rectangle of its buffer (segment 1's two
+ dimensions), the row dimension selected with a stride of 2 (rows
+ 1, 3 and 5). */
+
+ for (i = 0; i < DIM1; i++)
+ for (j = 0; j < DIM2; j++)
+ {
+#pragma omp target enter data map(to: x[i][j][:ROWS*COLS])
+ }
+
+ for (r = 0; r < ROWS; r++)
+ for (c = 0; c < COLS; c++)
+ x[1][1][r * COLS + c] = 100 * r + c;
+
+#pragma omp target update to((([ROWS][COLS]) x[1][1])[1:3:2][0:2])
+
+#pragma omp target map(alloc: x)
+ {
+ for (int r = 1; r <= 5; r += 2)
+ for (int c = 0; c < 2; c++)
+ x[1][1][r * COLS + c] += 1000;
+ }
+
+ /* Erase the host copies before pulling back from the device, so the
+ assertions below only see what actually made it across. */
+ for (i = 0; i < DIM1; i++)
+ for (j = 0; j < DIM2; j++)
+ memset (x[i][j], 0, ROWS * COLS * sizeof (int));
+
+ for (i = 0; i < DIM1; i++)
+ for (j = 0; j < DIM2; j++)
+ {
+#pragma omp target exit data map(from: x[i][j][:ROWS*COLS])
+ }
+
+ for (i = 0; i < DIM1; i++)
+ for (j = 0; j < DIM2; j++)
+ for (r = 0; r < ROWS; r++)
+ for (c = 0; c < COLS; c++)
+ {
+ int expected = 0;
+ if (i == 1 && j == 1 && (r == 1 || r == 3 || r == 5) && c < 2)
+ expected = 100 * r + c + 1000;
+ assert (x[i][j][r * COLS + c] == expected);
+ }
+
+ /* Case 2: update from device, same two-segment/two-dimension shape, on
+ a different pointer: clobber the host copy after pushing data to the
+ device, then confirm "update from" actually pulls the device's data
+ back rather than leaving the clobbered value. */
+
+ for (i = 0; i < DIM1; i++)
+ for (j = 0; j < DIM2; j++)
+ memset (x[i][j], 0, ROWS * COLS * sizeof (int));
+
+ for (i = 0; i < DIM1; i++)
+ for (j = 0; j < DIM2; j++)
+ {
+#pragma omp target enter data map(to: x[i][j][:ROWS*COLS])
+ }
+
+ for (r = 0; r < ROWS; r++)
+ for (c = 0; c < COLS; c++)
+ x[2][0][r * COLS + c] = 300 + 10 * r + c;
+
+#pragma omp target update to((([ROWS][COLS]) x[2][0:1])[2:2][1:3])
+
+#pragma omp target map(alloc: x)
+ {
+ for (int r = 2; r <= 3; r++)
+ for (int c = 1; c <= 3; c++)
+ x[2][0][r * COLS + c] += 1000;
+ }
+
+ for (r = 0; r < ROWS; r++)
+ for (c = 0; c < COLS; c++)
+ x[2][0][r * COLS + c] = -1;
+
+#pragma omp target update from((([ROWS][COLS]) x[2][0:1])[2:2][1:3])
+
+ for (i = 0; i < DIM1; i++)
+ for (j = 0; j < DIM2; j++)
+ for (r = 0; r < ROWS; r++)
+ for (c = 0; c < COLS; c++)
+ {
+ if (i == 2 && j == 0 && r >= 2 && r <= 3 && c >= 1 && c <= 3)
+ assert (x[i][j][r * COLS + c] == 300 + 10 * r + c + 1000);
+ else if (i == 2 && j == 0)
+ assert (x[i][j][r * COLS + c] == -1);
+ else
+ assert (x[i][j][r * COLS + c] == 0);
+ }
+
+ for (i = 0; i < DIM1; i++)
+ for (j = 0; j < DIM2; j++)
+ {
+#pragma omp target exit data map(delete: x[i][j][:ROWS*COLS])
+ }
+
+#pragma omp target exit data map(release: x[0:DIM1])
+
+ for (i = 0; i < DIM1; i++)
+ for (j = 0; j < DIM2; j++)
+ free (x[i][j]);
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/array-section-4.c b/libgomp/testsuite/libgomp.c-c++-common/array-section-4.c
new file mode 100644
index 00000000000..39afb53e216
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/array-section-4.c
@@ -0,0 +1,134 @@
+/* { dg-do run { target offload_device_nonshared_as } } */
+
+/* Test noncontiguous 'target update" through one level of pointer indirection,
+ where the array-of-pointers side has several dimensions selected by fixed
+ indices. Also check the transferred data can be accessed from device
+ code. */
+
+#include <string.h>
+#include <assert.h>
+#include <stdlib.h>
+
+#define D1 3
+#define D2 3
+#define D3 4
+#define M 8
+
+int main()
+{
+ int *w[D1][D2][D3];
+ int i, j, k, m;
+
+ for (i = 0; i < D1; i++)
+ for (j = 0; j < D2; j++)
+ for (k = 0; k < D3; k++)
+ w[i][j][k] = (int *) calloc (M, sizeof (int));
+
+#pragma omp target enter data map(alloc: w[0:D1])
+
+ /* Case 1: update to device. Three consecutive fixed indices select one
+ specific pointer, then a partial range of its buffer is updated. */
+
+ for (i = 0; i < D1; i++)
+ for (j = 0; j < D2; j++)
+ for (k = 0; k < D3; k++)
+ {
+#pragma omp target enter data map(to: w[i][j][k][:M])
+ }
+
+ for (m = 0; m < M; m++)
+ w[1][2][0][m] = 10 + m;
+
+#pragma omp target update to(w[1][2][0][2:4])
+
+#pragma omp target map(alloc: w)
+ {
+ for (int m = 2; m < 6; m++)
+ w[1][2][0][m] += 1000;
+ }
+
+ /* Erase the host copies before pulling back from the device, so the
+ assertions below only see what actually made it across. */
+ for (i = 0; i < D1; i++)
+ for (j = 0; j < D2; j++)
+ for (k = 0; k < D3; k++)
+ memset (w[i][j][k], 0, M * sizeof (int));
+
+ for (i = 0; i < D1; i++)
+ for (j = 0; j < D2; j++)
+ for (k = 0; k < D3; k++)
+ {
+#pragma omp target exit data map(from: w[i][j][k][:M])
+ }
+
+ for (i = 0; i < D1; i++)
+ for (j = 0; j < D2; j++)
+ for (k = 0; k < D3; k++)
+ for (m = 0; m < M; m++)
+ if (i == 1 && j == 2 && k == 0 && m >= 2 && m < 6)
+ assert (w[i][j][k][m] == 10 + m + 1000);
+ else
+ assert (w[i][j][k][m] == 0);
+
+ /* Case 2: update from device, same shape (three consecutive fixed indices),
+ through a different pointer: clobber the host copy after pushing data
+ to the device, then confirm "update from" actually pulls the
+ device's data back rather than leaving the clobbered value. */
+
+ for (i = 0; i < D1; i++)
+ for (j = 0; j < D2; j++)
+ for (k = 0; k < D3; k++)
+ memset (w[i][j][k], 0, M * sizeof (int));
+
+ for (i = 0; i < D1; i++)
+ for (j = 0; j < D2; j++)
+ for (k = 0; k < D3; k++)
+ {
+#pragma omp target enter data map(to: w[i][j][k][:M])
+ }
+
+ for (m = 0; m < M; m++)
+ w[2][0][3][m] = 100 + m;
+
+#pragma omp target update to(w[2][0][3][1:5])
+
+#pragma omp target map(alloc: w)
+ {
+ for (int m = 1; m < 6; m++)
+ w[2][0][3][m] += 1000;
+ }
+
+ for (m = 0; m < M; m++)
+ w[2][0][3][m] = -1;
+
+#pragma omp target update from(w[2][0][3][1:5])
+
+ for (i = 0; i < D1; i++)
+ for (j = 0; j < D2; j++)
+ for (k = 0; k < D3; k++)
+ for (m = 0; m < M; m++)
+ {
+ if (i == 2 && j == 0 && k == 3 && m >= 1 && m < 6)
+ assert (w[i][j][k][m] == 100 + m + 1000);
+ else if (i == 2 && j == 0 && k == 3)
+ assert (w[i][j][k][m] == -1);
+ else
+ assert (w[i][j][k][m] == 0);
+ }
+
+ for (i = 0; i < D1; i++)
+ for (j = 0; j < D2; j++)
+ for (k = 0; k < D3; k++)
+ {
+#pragma omp target exit data map(delete: w[i][j][k][:M])
+ }
+
+#pragma omp target exit data map(release: w[0:D1])
+
+ for (i = 0; i < D1; i++)
+ for (j = 0; j < D2; j++)
+ for (k = 0; k < D3; k++)
+ free (w[i][j][k]);
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/array-section-5.c b/libgomp/testsuite/libgomp.c-c++-common/array-section-5.c
new file mode 100644
index 00000000000..63d798aaa7c
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/array-section-5.c
@@ -0,0 +1,136 @@
+/* { dg-do run { target offload_device_nonshared_as } } */
+
+/* Test noncontiguous target update through one level of pointer indirection,
+ where the array-shaping cast has more dimensions than are explicitly
+ sectioned afterwards. Also check that transferred data can be read and
+ written from device code. */
+
+#include <string.h>
+#include <assert.h>
+#include <stdlib.h>
+
+#define DIM1 3
+#define DIM2 4
+#define ROWS 6
+#define COLS 6
+
+int main()
+{
+ int *x[DIM1][DIM2];
+ int i, j, r, c;
+
+ for (i = 0; i < DIM1; i++)
+ for (j = 0; j < DIM2; j++)
+ x[i][j] = (int *) calloc (ROWS * COLS, sizeof (int));
+
+#pragma omp target enter data map(alloc: x[0:DIM1])
+
+ /* Case 1: update to device. Two fixed indices select a single pointer
+ from the 2D array-of-pointers (segment 0); the selected pointer
+ contributes a strided range of whole rows (segment 1's only explicit
+ dimension). */
+
+ for (i = 0; i < DIM1; i++)
+ for (j = 0; j < DIM2; j++)
+ {
+#pragma omp target enter data map(to: x[i][j][:ROWS*COLS])
+ }
+
+ for (r = 0; r < ROWS; r++)
+ for (c = 0; c < COLS; c++)
+ x[1][1][r * COLS + c] = 100 * r + c;
+
+#pragma omp target update to((([ROWS][COLS]) x[1][1])[1:3:2])
+
+#pragma omp target map(alloc: x)
+ {
+ for (int r = 1; r <= 5; r += 2)
+ for (int c = 0; c < COLS; c++)
+ x[1][1][r * COLS + c] += 1000;
+ }
+
+ /* Erase the host copies before pulling back from the device, so the
+ assertions below only see what actually made it across. */
+ for (i = 0; i < DIM1; i++)
+ for (j = 0; j < DIM2; j++)
+ memset (x[i][j], 0, ROWS * COLS * sizeof (int));
+
+ for (i = 0; i < DIM1; i++)
+ for (j = 0; j < DIM2; j++)
+ {
+#pragma omp target exit data map(from: x[i][j][:ROWS*COLS])
+ }
+
+ for (i = 0; i < DIM1; i++)
+ for (j = 0; j < DIM2; j++)
+ for (r = 0; r < ROWS; r++)
+ for (c = 0; c < COLS; c++)
+ {
+ int expected = 0;
+ if (i == 1 && j == 1 && (r == 1 || r == 3 || r == 5))
+ expected = 100 * r + c + 1000;
+ assert (x[i][j][r * COLS + c] == expected);
+ }
+
+ /* Case 2: update from device, same shape (rows selected, columns left
+ to their whole span), on a different pointer: clobber the host copy
+ after pushing data to the device, then confirm "update from" actually
+ pulls the device's data back rather than leaving the clobbered
+ value. */
+
+ for (i = 0; i < DIM1; i++)
+ for (j = 0; j < DIM2; j++)
+ memset (x[i][j], 0, ROWS * COLS * sizeof (int));
+
+ for (i = 0; i < DIM1; i++)
+ for (j = 0; j < DIM2; j++)
+ {
+#pragma omp target enter data map(to: x[i][j][:ROWS*COLS])
+ }
+
+ for (r = 0; r < ROWS; r++)
+ for (c = 0; c < COLS; c++)
+ x[2][0][r * COLS + c] = 300 + 10 * r + c;
+
+#pragma omp target update to((([ROWS][COLS]) x[2][0:1])[2:2])
+
+#pragma omp target map(alloc: x)
+ {
+ for (int r = 2; r <= 3; r++)
+ for (int c = 0; c < COLS; c++)
+ x[2][0][r * COLS + c] += 1000;
+ }
+
+ for (r = 0; r < ROWS; r++)
+ for (c = 0; c < COLS; c++)
+ x[2][0][r * COLS + c] = -1;
+
+#pragma omp target update from((([ROWS][COLS]) x[2][0:1])[2:2])
+
+ for (i = 0; i < DIM1; i++)
+ for (j = 0; j < DIM2; j++)
+ for (r = 0; r < ROWS; r++)
+ for (c = 0; c < COLS; c++)
+ {
+ if (i == 2 && j == 0 && r >= 2 && r <= 3)
+ assert (x[i][j][r * COLS + c] == 300 + 10 * r + c + 1000);
+ else if (i == 2 && j == 0)
+ assert (x[i][j][r * COLS + c] == -1);
+ else
+ assert (x[i][j][r * COLS + c] == 0);
+ }
+
+ for (i = 0; i < DIM1; i++)
+ for (j = 0; j < DIM2; j++)
+ {
+#pragma omp target exit data map(delete: x[i][j][:ROWS*COLS])
+ }
+
+#pragma omp target exit data map(release: x[0:DIM1])
+
+ for (i = 0; i < DIM1; i++)
+ for (j = 0; j < DIM2; j++)
+ free (x[i][j]);
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/array-section-6.c b/libgomp/testsuite/libgomp.c-c++-common/array-section-6.c
new file mode 100644
index 00000000000..01127006d55
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/array-section-6.c
@@ -0,0 +1,75 @@
+/* { dg-do run { target offload_device_nonshared_as } } */
+
+/* Test array-shaping cast applied to a plain pointer, with a further array
+ section applied on top of the cast. */
+
+#include <string.h>
+#include <assert.h>
+#include <stdlib.h>
+
+#define N 100
+
+int main()
+{
+ int *ptr = (int *) calloc (N, sizeof (int));
+ int i;
+
+#pragma omp target enter data map(to: ptr[:N])
+
+ for (i = 0; i < N; i++)
+ ptr[i] = i;
+
+#pragma omp target update to((([N]) ptr)[10:30])
+
+#pragma omp target map(alloc: ptr[:N])
+ {
+ for (int i = 10; i < 40; i++)
+ ptr[i] += 1000;
+ }
+
+ /* Erase the host copy before pulling back from the device, so the
+ assertions below only see what actually made it across. */
+ memset (ptr, 0, N * sizeof (int));
+
+#pragma omp target exit data map(from: ptr[:N])
+
+ for (i = 0; i < N; i++)
+ if (i >= 10 && i < 40)
+ assert (ptr[i] == i + 1000);
+ else
+ assert (ptr[i] == 0);
+
+ /* Update from device: clobber the host copy after pushing data to the
+ device, then confirm "update from" actually pulls the device's data
+ back rather than leaving the clobbered value. */
+
+#pragma omp target enter data map(to: ptr[:N])
+
+ for (i = 0; i < N; i++)
+ ptr[i] = i;
+
+#pragma omp target update to((([N]) ptr)[10:30])
+
+#pragma omp target map(alloc: ptr[:N])
+ {
+ for (int i = 10; i < 40; i++)
+ ptr[i] += 1000;
+ }
+
+ for (i = 0; i < N; i++)
+ ptr[i] = -1;
+
+#pragma omp target update from((([N]) ptr)[10:30])
+
+ for (i = 0; i < N; i++)
+ if (i >= 10 && i < 40)
+ assert (ptr[i] == i + 1000);
+ else
+ assert (ptr[i] == -1);
+
+#pragma omp target exit data map(delete: ptr[:N])
+
+ free (ptr);
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/array-section-7.c b/libgomp/testsuite/libgomp.c-c++-common/array-section-7.c
new file mode 100644
index 00000000000..b75ccea2539
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/array-section-7.c
@@ -0,0 +1,104 @@
+/* { dg-do run { target offload_device_nonshared_as } } */
+
+/* Test array-shaping cast applied directly to a fixed-index element of a 1D
+ array-of-pointers, with no further outer section.. */
+
+#include <string.h>
+#include <assert.h>
+#include <stdlib.h>
+
+#define DIM1 5
+#define N 20
+
+int main ()
+{
+ int *x[DIM1];
+ int i, k;
+
+ for (i = 0; i < DIM1; i++)
+ x[i] = (int *) calloc (N, sizeof (int));
+
+#pragma omp target enter data map(alloc: x[0:DIM1])
+
+ for (i = 0; i < DIM1; i++)
+ {
+#pragma omp target enter data map(to: x[i][:N])
+ }
+
+ /* Case 1: update to device, single fixed pointer index, whole span. */
+
+ for (k = 0; k < N; k++)
+ x[2][k] = 100 + k;
+
+#pragma omp target update to(([N]) x[2])
+
+#pragma omp target map(alloc: x)
+ {
+ for (int k = 0; k < N; k++)
+ x[2][k] += 1000;
+ }
+
+ for (i = 0; i < DIM1; i++)
+ memset (x[i], 0, N * sizeof (int));
+
+ for (i = 0; i < DIM1; i++)
+ {
+#pragma omp target exit data map(from: x[i][:N])
+ }
+
+ for (i = 0; i < DIM1; i++)
+ for (k = 0; k < N; k++)
+ {
+ int expected = (i == 2) ? 100 + k + 1000 : 0;
+ assert (x[i][k] == expected);
+ }
+
+ /* Case 2: update from device, same shape, different index, confirming
+ "update from" actually pulls the device's data back. */
+
+ for (i = 0; i < DIM1; i++)
+ {
+#pragma omp target enter data map(to: x[i][:N])
+ }
+
+ for (k = 0; k < N; k++)
+ x[3][k] = 200 + k;
+
+#pragma omp target update to(([N]) x[3])
+
+#pragma omp target map(alloc: x)
+ {
+ for (int k = 0; k < N; k++)
+ x[3][k] += 1000;
+ }
+
+ for (k = 0; k < N; k++)
+ x[3][k] = -1;
+
+#pragma omp target update from(([N]) x[3])
+
+ for (i = 0; i < DIM1; i++)
+ for (k = 0; k < N; k++)
+ {
+ int expected;
+ if (i == 3)
+ expected = 200 + k + 1000;
+ else if (i == 2)
+ expected = 100 + k + 1000;
+ else
+ expected = 0;
+ assert (x[i][k] == expected);
+ }
+
+ for (i = 0; i < DIM1; i++)
+ {
+#pragma omp target exit data map(delete: x[i][:N])
+ }
+
+#pragma omp target exit data map(release: x[0:DIM1])
+
+ for (i = 0; i < DIM1; i++)
+ free (x[i]);
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/array-section-8.c b/libgomp/testsuite/libgomp.c-c++-common/array-section-8.c
new file mode 100644
index 00000000000..eccc187dbd6
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/array-section-8.c
@@ -0,0 +1,87 @@
+/* { dg-do run { target offload_device_nonshared_as } } */
+
+/* Test three-segment (two pointer indirections) noncontiguous target update:
+ a strided section reached by crossing two levels of pointer indirection.
+ Only one dimension per segment. */
+
+#include <assert.h>
+#include <stdlib.h>
+
+#define DIM1 4
+#define DIM2 3
+#define N 20
+
+int
+main ()
+{
+ int **arr[DIM1];
+ int i, j, k;
+
+ for (i = 0; i < DIM1; i++)
+ {
+ arr[i] = (int **) malloc (DIM2 * sizeof (int *));
+ for (j = 0; j < DIM2; j++)
+ arr[i][j] = (int *) calloc (N, sizeof (int));
+ }
+
+ for (i = 0; i < DIM1; i++)
+ for (j = 0; j < DIM2; j++)
+ for (k = 0; k < N; k++)
+ arr[i][j][k] = i * 10000 + j * 100 + k;
+
+ for (i = 0; i < DIM1; i++)
+ for (j = 0; j < DIM2; j++)
+ {
+#pragma omp target enter data map(to: arr[i][j][:N])
+ }
+
+ /* Corrupt the host copy of exactly the two rows the update below will
+ touch, so matching the original values afterward proves the update
+ actually pulled from the device instead of being a no-op. */
+ for (j = 0; j < DIM2; j += 2)
+ for (k = 0; k < N; k++)
+ arr[2][j][k] = -1;
+
+#pragma omp target update from(arr[2][0:2:2][3:5:2])
+
+ {
+ int touched[5] = { 3, 5, 7, 9, 11 };
+
+ for (i = 0; i < DIM1; i++)
+ for (j = 0; j < DIM2; j++)
+ for (k = 0; k < N; k++)
+ {
+ int expected;
+ if (i == 2 && (j == 0 || j == 2))
+ {
+ int is_touched = 0, t;
+ for (t = 0; t < 5; t++)
+ if (touched[t] == k)
+ is_touched = 1;
+ /* Restored from the device if touched, still corrupted
+ otherwise -- proves the update copied exactly the
+ requested elements. */
+ expected = is_touched ? i * 10000 + j * 100 + k : -1;
+ }
+ else
+ /* Untouched by either the corruption or the update. */
+ expected = i * 10000 + j * 100 + k;
+ assert (arr[i][j][k] == expected);
+ }
+ }
+
+ for (i = 0; i < DIM1; i++)
+ for (j = 0; j < DIM2; j++)
+ {
+#pragma omp target exit data map(delete: arr[i][j][:N])
+ }
+
+ for (i = 0; i < DIM1; i++)
+ {
+ for (j = 0; j < DIM2; j++)
+ free (arr[i][j]);
+ free (arr[i]);
+ }
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/array-section-9.c b/libgomp/testsuite/libgomp.c-c++-common/array-section-9.c
new file mode 100644
index 00000000000..6fc152dcd1f
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/array-section-9.c
@@ -0,0 +1,120 @@
+/* { dg-do run { target offload_device_nonshared_as } } */
+
+/* Test three-segment (two pointer indirections) noncontiguous target update,
+ with *two* dimensions selected per segment -- six dimensions total. */
+
+#include <assert.h>
+#include <stdlib.h>
+
+#define DIM1A 2
+#define DIM1B 3
+#define ROWS1 3
+#define NCOLS 3
+#define ROWS2 3
+#define ROWLEN 20
+
+typedef int (*p2_t)[ROWLEN];
+typedef p2_t (*p1_t)[NCOLS];
+
+static int
+encode (int i, int j, int r1, int c1, int r2, int k)
+{
+ return i * 1000000 + j * 100000 + r1 * 10000 + c1 * 1000 + r2 * 100 + k;
+}
+
+int
+main ()
+{
+ p1_t arr[DIM1A][DIM1B];
+ int i, j, r1, c1, r2, k;
+
+ for (i = 0; i < DIM1A; i++)
+ for (j = 0; j < DIM1B; j++)
+ {
+ arr[i][j] = (p1_t) malloc (ROWS1 * NCOLS * sizeof (p2_t));
+ for (r1 = 0; r1 < ROWS1; r1++)
+ for (c1 = 0; c1 < NCOLS; c1++)
+ arr[i][j][r1][c1]
+ = (p2_t) malloc (ROWS2 * ROWLEN * sizeof (int));
+ }
+
+ for (i = 0; i < DIM1A; i++)
+ for (j = 0; j < DIM1B; j++)
+ for (r1 = 0; r1 < ROWS1; r1++)
+ for (c1 = 0; c1 < NCOLS; c1++)
+ for (r2 = 0; r2 < ROWS2; r2++)
+ for (k = 0; k < ROWLEN; k++)
+ arr[i][j][r1][c1][r2][k] = encode (i, j, r1, c1, r2, k);
+
+ for (i = 0; i < DIM1A; i++)
+ for (j = 0; j < DIM1B; j++)
+ for (r1 = 0; r1 < ROWS1; r1++)
+ for (c1 = 0; c1 < NCOLS; c1++)
+ {
+#pragma omp target enter data map(to: arr[i][j][r1][c1][:ROWS2])
+ }
+
+ /* Corrupt the whole of ARR[1]'s data (every J/R1/C1/R2/K), so that
+ precisely matching the six-dimensional touched set below after the
+ update proves each dimension of each segment was walked correctly. */
+ for (j = 0; j < DIM1B; j++)
+ for (r1 = 0; r1 < ROWS1; r1++)
+ for (c1 = 0; c1 < NCOLS; c1++)
+ for (r2 = 0; r2 < ROWS2; r2++)
+ for (k = 0; k < ROWLEN; k++)
+ arr[1][j][r1][c1][r2][k] = -1;
+
+#pragma omp target update from(arr[1][0:2:2][0:2:2][0:2:2][0:2][3:5:2])
+
+ for (i = 0; i < DIM1A; i++)
+ for (j = 0; j < DIM1B; j++)
+ for (r1 = 0; r1 < ROWS1; r1++)
+ for (c1 = 0; c1 < NCOLS; c1++)
+ for (r2 = 0; r2 < ROWS2; r2++)
+ for (k = 0; k < ROWLEN; k++)
+ {
+ int expected;
+ if (i == 1)
+ {
+ int j_touched = (j == 0 || j == 2);
+ int r1_touched = (r1 == 0 || r1 == 2);
+ int c1_touched = (c1 == 0 || c1 == 2);
+ int r2_touched = (r2 == 0 || r2 == 1);
+ int k_touched
+ = (k == 3 || k == 5 || k == 7 || k == 9 || k == 11);
+ if (j_touched && r1_touched && c1_touched
+ && r2_touched && k_touched)
+ /* Restored from the device -- proves this exact
+ combination of dimensions across all three
+ segments was selected. */
+ expected = encode (i, j, r1, c1, r2, k);
+ else
+ /* Still corrupted -- proves the update did not
+ touch anything outside the requested set. */
+ expected = -1;
+ }
+ else
+ /* Untouched by either the corruption or the update. */
+ expected = encode (i, j, r1, c1, r2, k);
+ assert (arr[i][j][r1][c1][r2][k] == expected);
+ }
+
+ for (i = 0; i < DIM1A; i++)
+ for (j = 0; j < DIM1B; j++)
+ for (r1 = 0; r1 < ROWS1; r1++)
+ for (c1 = 0; c1 < NCOLS; c1++)
+ {
+#pragma omp target exit data map(delete: arr[i][j][r1][c1][:ROWS2])
+ }
+
+ for (i = 0; i < DIM1A; i++)
+ for (j = 0; j < DIM1B; j++)
+ {
+ for (r1 = 0; r1 < ROWS1; r1++)
+ for (c1 = 0; c1 < NCOLS; c1++)
+ free (arr[i][j][r1][c1]);
+ free (arr[i][j]);
+ }
+
+ return 0;
+}
--
2.53.0