[gcc/devel/omp/gcc-14] OpenMP: Dimension ordering for array-shaping operator for C and C++

[Paul-Antoine Arras]

https://gcc.gnu.org/g:51eafa3528f79e8817da82892358ad10a4c65f08

commit 51eafa3528f79e8817da82892358ad10a4c65f08
Author: Julian Brown <julian@codesourcery.com>
Date:   Thu Jul 13 10:32:16 2023 +0000

    OpenMP: Dimension ordering for array-shaping operator for C and C++
    
    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.
    
    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.

Diff:
---
 gcc/ChangeLog.omp                                  |  5 ++++
 gcc/fortran/ChangeLog.omp                          |  5 ++++
 gcc/fortran/trans-openmp.cc                        |  2 +-
 gcc/omp-low.cc                                     |  6 ++--
 libgomp/ChangeLog.omp                              |  4 +++
 .../libgomp.c-c++-common/array-shaping-14.c        | 34 ++++++++++++++++++++++
 6 files changed, 52 insertions(+), 4 deletions(-)

diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp
index 1b2266c5c0e..71b160a3941 100644
--- a/gcc/ChangeLog.omp
+++ b/gcc/ChangeLog.omp
@@ -1,3 +1,8 @@
+2023-07-14  Julian Brown  <julian@codesourcery.com>
+
+	* omp-low.cc (lower_omp_target): Reverse iteration over array
+	dimensions.
+
 2023-07-12  Julian Brown  <julian@codesourcery.com>
 
 	* omp-low.cc (lower_omp_target): Calculate volume enclosing
diff --git a/gcc/fortran/ChangeLog.omp b/gcc/fortran/ChangeLog.omp
index 3794560b606..a60aafb0f5e 100644
--- a/gcc/fortran/ChangeLog.omp
+++ b/gcc/fortran/ChangeLog.omp
@@ -1,3 +1,8 @@
+2023-07-14  Julian Brown  <julian@codesourcery.com>
+
+	* trans-openmp.cc (gfc_trans_omp_arrayshape_type): Reverse dimension
+	ordering for created array type.
+
 2023-07-03  Julian Brown  <julian@codesourcery.com>
 
 	* trans-openmp.cc (gfc_omp_deep_map_kind_p): Handle
diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc
index e4475cfc403..e5cd653c536 100644
--- a/gcc/fortran/trans-openmp.cc
+++ b/gcc/fortran/trans-openmp.cc
@@ -4261,7 +4261,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.  */
diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc
index 0caf5dea05f..37be9d46b3e 100644
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -13844,7 +13844,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;
@@ -13919,7 +13919,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
@@ -13957,7 +13957,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;
 		      }
diff --git a/libgomp/ChangeLog.omp b/libgomp/ChangeLog.omp
index e00cb42ef53..3a4a9c76de4 100644
--- a/libgomp/ChangeLog.omp
+++ b/libgomp/ChangeLog.omp
@@ -1,3 +1,7 @@
+2023-07-14  Julian Brown  <julian@codesourcery.com>
+
+	* testsuite/libgomp.c-c++-common/array-shaping-14.c: New test.
+
 2023-09-05  Julian Brown  <julian@codesourcery.com>
 
 	* testsuite/libgomp.c/array-shaping-1.c: New test.
diff --git a/libgomp/testsuite/libgomp.c-c++-common/array-shaping-14.c b/libgomp/testsuite/libgomp.c-c++-common/array-shaping-14.c
new file mode 100644
index 00000000000..4ca6f794f93
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/array-shaping-14.c
@@ -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;
+}