[og13] OpenMP: Dimension ordering for array-shaping operator for C and C++
Checks
Commit Message
This patch fixes a bug in non-contiguous 'target update' operations using
the new array-shaping operator for C and C++, processing the dimensions
of the array the wrong way round during the OpenMP lowering pass.
Fortran was also incorrectly using the wrong ordering but the second
reversal in omp-low.cc made it produce the correct result.
The C and C++ bug only affected array shapes where the dimension sizes
are different ([X][Y]) - several existing tests used the same value
for both/all dimensions ([X][X]), which masked the problem. Only the
array dimensions (extents) are affected, not e.g. the indices, lengths
or strides for array sections.
This patch reverses the order used in both omp-low.cc and the Fortran
front-end, so the order should now be correct for all supported base
languages.
Tested with offloading to AMD GCN. I will apply (to og13) shortly.
2023-07-14 Julian Brown <julian@codesourcery.com>
gcc/fortran/
* trans-openmp.cc (gfc_trans_omp_arrayshape_type): Reverse dimension
ordering for created array type.
gcc/
* omp-low.cc (lower_omp_target): Reverse iteration over array
dimensions.
libgomp/
* testsuite/libgomp.c-c++-common/array-shaping-14.c: New test.
---
gcc/fortran/trans-openmp.cc | 2 +-
gcc/omp-low.cc | 6 ++--
.../libgomp.c-c++-common/array-shaping-14.c | 34 +++++++++++++++++++
3 files changed, 38 insertions(+), 4 deletions(-)
create mode 100644 libgomp/testsuite/libgomp.c-c++-common/array-shaping-14.c
@@ -4271,7 +4271,7 @@ gfc_trans_omp_arrayshape_type (tree type, vec<tree> *dims)
{
gcc_assert (dims->length () > 0);
- for (int i = dims->length () - 1; i >= 0; i--)
+ for (unsigned i = 0; i < dims->length (); i++)
{
tree dim = fold_convert (sizetype, (*dims)[i]);
/* We need the index of the last element, not the array size. */
@@ -14290,7 +14290,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
dims++;
}
- int tdim = tdims.length () - 1;
+ unsigned tdim = 0;
vec<constructor_elt, va_gc> *vdim;
vec<constructor_elt, va_gc> *vindex;
@@ -14365,7 +14365,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
nc = nc2;
}
- if (tdim >= 0)
+ if (tdim < tdims.length ())
{
/* We have an array shape -- use that to find the
total size of the data on the target to look up
@@ -14403,7 +14403,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
"for array");
dim = index = len = stride = error_mark_node;
}
- tdim--;
+ tdim++;
c = nc;
}
new file mode 100644
@@ -0,0 +1,34 @@
+/* { dg-do run { target offload_device_nonshared_as } } */
+
+#include <stdlib.h>
+#include <stdio.h>
+#include <assert.h>
+
+typedef struct {
+ int *ptr;
+} S;
+
+int main(void)
+{
+ S q;
+ q.ptr = (int *) calloc (9 * 11, sizeof (int));
+
+#pragma omp target enter data map(to: q.ptr, q.ptr[0:9*11])
+
+#pragma omp target
+ for (int i = 0; i < 9*11; i++)
+ q.ptr[i] = i;
+
+#pragma omp target update from(([9][11]) q.ptr[3:3:2][1:4:3])
+
+ for (int j = 0; j < 9; j++)
+ for (int i = 0; i < 11; i++)
+ if (j >= 3 && j <= 7 && ((j - 3) % 2) == 0
+ && i >= 1 && i <= 10 && ((i - 1) % 3) == 0)
+ assert (q.ptr[j * 11 + i] == j * 11 + i);
+ else
+ assert (q.ptr[j * 11 + i] == 0);
+
+#pragma omp target exit data map(release: q.ptr, q.ptr[0:9*11])
+ return 0;
+}