[gcc r12-9841] omp-expand.cc: Fix wrong code with non-rectangular loop nest [PR111017]

[gcc r12-9841] omp-expand.cc: Fix wrong code with non-rectangular loop nest [PR111017]

[Tobias Burnus]

https://gcc.gnu.org/g:2f8ccacd41a45bccf3e19625c0fbd2627472b45e

commit r12-9841-g2f8ccacd41a45bccf3e19625c0fbd2627472b45e
Author: Tobias Burnus <tobias@codesourcery.com>
Date:   Sat Aug 19 07:49:06 2023 +0200

    omp-expand.cc: Fix wrong code with non-rectangular loop nest [PR111017]
    
    Before commit r12-5295-g47de0b56ee455e, all gimple_build_cond in
    expand_omp_for_* were inserted with
      gsi_insert_before (gsi_p, cond_stmt, GSI_SAME_STMT);
    except the one dealing with the multiplicative factor that was
      gsi_insert_after (gsi, cond_stmt, GSI_CONTINUE_LINKING);
    
    That commit for PR103208 fixed the issue of some missing regimplify of
    operands of GIMPLE_CONDs by moving the condition handling to the new function
    expand_omp_build_cond. While that function has an 'bool after = false'
    argument to switch between the two variants.
    
    However, all callers ommited this argument. This commit reinstates the
    prior behavior by passing 'true' for the factor != 0 condition, fixing
    the included testcase.
    
            PR middle-end/111017
    gcc/
            * omp-expand.cc (expand_omp_for_init_vars): Pass after=true
            to expand_omp_build_cond for 'factor != 0' condition, resulting
            in pre-r12-5295-g47de0b56ee455e code for the gimple insert.
    
    libgomp/
            * testsuite/libgomp.c-c++-common/non-rect-loop-1.c: New test.
    
    (cherry picked from commit 1dc65003b66e5a97200f454eeddcccfce34416b3)

Diff:
---
 gcc/omp-expand.cc                                  |  3 +-
 .../libgomp.c-c++-common/non-rect-loop-1.c         | 72 ++++++++++++++++++++++
 2 files changed, 74 insertions(+), 1 deletion(-)

diff --git a/gcc/omp-expand.cc b/gcc/omp-expand.cc
index c3d6447440d2..09190a79ab2e 100644
--- a/gcc/omp-expand.cc
+++ b/gcc/omp-expand.cc
@@ -2564,7 +2564,8 @@ expand_omp_for_init_vars (struct omp_for_data *fd, gimple_stmt_iterator *gsi,
 	      tree factor = fd->factor;
 	      gcond *cond_stmt
 		= expand_omp_build_cond (gsi, NE_EXPR, factor,
-					 build_zero_cst (TREE_TYPE (factor)));
+					 build_zero_cst (TREE_TYPE (factor)),
+					 true);
 	      edge e = split_block (gsi_bb (*gsi), cond_stmt);
 	      basic_block bb0 = e->src;
 	      e->flags = EDGE_TRUE_VALUE;
diff --git a/libgomp/testsuite/libgomp.c-c++-common/non-rect-loop-1.c b/libgomp/testsuite/libgomp.c-c++-common/non-rect-loop-1.c
new file mode 100644
index 000000000000..fbd462b36831
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/non-rect-loop-1.c
@@ -0,0 +1,72 @@
+/* PR middle-end/111017  */
+
+#include <omp.h>
+
+#define DIM 32
+#define N (DIM*DIM)
+
+int
+main ()
+{
+  int a[N], b[N], c[N];
+  int dim = DIM;
+
+  for (int i = 0; i < N; i++)
+    {
+      a[i] = 3*i;
+      b[i] = 7*i;
+      c[i] = 42;
+    }
+
+  #pragma omp parallel for collapse(2)
+  for (int i = 0; i < DIM; i++)
+    for (int j = (i*DIM); j < (i*DIM + DIM); j++)
+      c[j] = a[j] + b[j];
+
+  for (int i = 0; i < DIM; i++)
+    for (int j = (i*DIM); j < (i*DIM + DIM); j++)
+      if (c[j] != a[j] + b[j] || c[j] != 3*j +7*j)
+	__builtin_abort ();
+  for (int i = 0; i < N; i++)
+    c[i] = 42;
+
+  #pragma omp parallel for collapse(2)
+  for (int i = 0; i < dim; i++)
+    for (int j = (i*dim); j < (i*dim + dim); j++)
+      c[j] = a[j] + b[j];
+
+  for (int i = 0; i < DIM; i++)
+    for (int j = (i*DIM); j < (i*DIM + DIM); j++)
+      if (c[j] != a[j] + b[j] || c[j] != 3*j +7*j)
+	__builtin_abort ();
+  for (int i = 0; i < N; i++)
+    c[i] = 42;
+
+  for (int dev = 0; dev <= omp_get_num_devices(); dev++)
+    {
+      #pragma omp target teams loop device(dev) map(to:a,b) map(from:c)
+      for (int i = 0; i < DIM; i++)
+	for (int j = (i*DIM); j < (i*DIM + DIM); j++)
+	  c[j] = a[j] + b[j];
+
+      for (int i = 0; i < DIM; i++)
+	for (int j = (i*DIM); j < (i*DIM + DIM); j++)
+	  if (c[j] != a[j] + b[j] || c[j] != 3*j +7*j)
+	    __builtin_abort ();
+      for (int i = 0; i < N; i++)
+	c[i] = 42;
+
+      #pragma omp target teams loop device(dev) map(to:a,b) map(from:c)
+      for (int i = 0; i < dim; i++)
+	for (int j = (i*dim); j < (i*dim + dim); j++)
+	  c[j] = a[j] + b[j];
+
+      for (int i = 0; i < DIM; i++)
+	for (int j = (i*DIM); j < (i*DIM + DIM); j++)
+	  if (c[j] != a[j] + b[j] || c[j] != 3*j +7*j)
+	    __builtin_abort ();
+      for (int i = 0; i < N; i++)
+	c[i] = 42;
+    }
+  return 0;
+}