[gcc r10-9711] [OpenACC] Fix an ICE where a loop with GT condition is collapsed.

[gcc r10-9711] [OpenACC] Fix an ICE where a loop with GT condition is collapsed.

[Hafiz Abid Qadeer]

https://gcc.gnu.org/g:e4dcb3383bff4c209a918127551cabc56b4171ae

commit r10-9711-ge4dcb3383bff4c209a918127551cabc56b4171ae
Author: Hafiz Abid Qadeer <abidh@codesourcery.com>
Date:   Thu Apr 8 17:31:30 2021 +0100

    [OpenACC] Fix an ICE where a loop with GT condition is collapsed.
    
    We have seen an ICE both on trunk and devel/omp/gcc-10 branches which can
    be reprodued with this simple testcase.  It occurs if an OpenACC loop has
    a collapse clause and any of the loop being collapsed uses GT or GE
    condition.  This issue is specific to OpenACC.
    
    int main (void)
    {
      int ix, iy;
      int dim_x = 16, dim_y = 16;
      {
           for (iy = dim_y - 1; iy > 0; --iy)
           for (ix = dim_x - 1; ix > 0; --ix)
            ;
      }
    }
    
    The problem is caused by a failing assertion in expand_oacc_collapse_init.
    It checks that cond_code for fd->loop should be same as cond_code for all
    the loops that are being collapsed.  As the cond_code for fd->loop is
    LT_EXPR with collapse clause (set at the end of omp_extract_for_data),
    this assertion forces that all the loop in collapse clause should use
    < operator.
    
    There does not seem to be anything in the code which demands this
    condition as loop with > condition works ok otherwise.  I digged old
    mailing list a bit but could not find any discussion on this change.
    Looking at the code, expand_oacc_for checks that fd->loop->cond_code is
    either LT_EXPR or GT_EXPR.  I guess the original intention was to have
    similar checks on the loop which are being collapsed. But the way check
    was written does not acheive that.
    
    I have fixed it by modifying the check in the assertion to be same as
    check on fd->loop->cond_code.
    
    I tested goacc and libgomp (with nvptx offloading) and did not see any
    regression.  I have added new tests to check collapse with GT/GE condition.
    
            PR middle-end/98088
            gcc/
            * omp-expand.c (expand_oacc_collapse_init): Update condition in
            a gcc_assert.
    
            gcc/testsuite/
            * c-c++-common/goacc/collapse-2.c: New.
    
            libgomp/
            * testsuite/libgomp.oacc-c-c++-common/collapse-2.c: Add check
            for loop with GT/GE condition.
            * testsuite/libgomp.oacc-c-c++-common/collapse-3.c: Likewise.
    
    (cherry picked from commit ac200799acb5cd2fb9e1758f6bf5fff1978daaeb)

Diff:
---
 gcc/omp-expand.c                                   |  2 +-
 gcc/testsuite/c-c++-common/goacc/collapse-2.c      | 56 ++++++++++++++++++++++
 .../libgomp.oacc-c-c++-common/collapse-2.c         | 17 ++++++-
 .../libgomp.oacc-c-c++-common/collapse-3.c         | 15 +++++-
 4 files changed, 85 insertions(+), 5 deletions(-)

diff --git a/gcc/omp-expand.c b/gcc/omp-expand.c
index 90df5c90547..b73fcf0b0a2 100644
--- a/gcc/omp-expand.c
+++ b/gcc/omp-expand.c
@@ -1569,7 +1569,7 @@ expand_oacc_collapse_init (const struct omp_for_data *fd,
       tree diff_type = iter_type;
       tree plus_type = iter_type;
 
-      gcc_assert (loop->cond_code == fd->loop.cond_code);
+      gcc_assert (loop->cond_code == LT_EXPR || loop->cond_code == GT_EXPR);
 
       if (POINTER_TYPE_P (iter_type))
 	plus_type = sizetype;
diff --git a/gcc/testsuite/c-c++-common/goacc/collapse-2.c b/gcc/testsuite/c-c++-common/goacc/collapse-2.c
new file mode 100644
index 00000000000..e46028cd5c4
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/collapse-2.c
@@ -0,0 +1,56 @@
+/* Test for ICE as reported in PR98088.  */
+
+int i, j;
+
+void
+f1 (void)
+{
+  #pragma acc parallel
+  #pragma acc loop collapse (2)
+  for (i = 5; i > 5; i--)
+	for (j = 5; j > 0; j--)
+	  ;
+}
+
+void
+f2 (void)
+{
+  #pragma acc parallel
+  #pragma acc loop collapse (2)
+  for (i = 0; i < 5; i++)
+	for (j = 5; j > 0; j--)
+	  ;
+}
+
+void
+f3 (void)
+{
+  #pragma acc parallel
+  #pragma acc loop collapse (2)
+  for (i = 5; i >= 0; i--)
+	for (j = 5; j >= 0; j--)
+	  ;
+}
+
+void f4 ()
+{
+  #pragma acc parallel loop tile(2, 3)
+  for (int i = 0; i > 8; i++)
+    for (int j = 0; j > 8; j++);
+}
+
+void f5 ()
+{
+  #pragma acc parallel loop tile(2, 3)
+  for (int i = 0; i > 8; i++)
+    for (long j = 0; j > 8; j++);
+}
+
+void
+f6 (int a[32][32])
+{
+  #pragma acc parallel loop collapse(2)
+  for (int i = 16; i > 8; i--)
+    for (int j = 16; j > 8; j--)
+      a[i][j] = i + j;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/collapse-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/collapse-2.c
index 1ea0a6b846d..7a8cfd2f3d4 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/collapse-2.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/collapse-2.c
@@ -5,7 +5,7 @@
 int
 main (void)
 {
-  int i, j, k, l = 0, f = 0, x = 0;
+  int i, j, k, l = 0, f = 0, x = 0, l2 = 0;
   int m1 = 4, m2 = -5, m3 = 17;
 
 #pragma acc parallel
@@ -20,6 +20,19 @@ main (void)
 	    }
 	}
 
+  /*  Test loop with > condition.  */
+#pragma acc parallel
+  #pragma acc loop seq collapse(3) reduction(+:l2)
+    for (i = -2; i < m1; i++)
+      for (j = -3; j > (m2 - 1); j--)
+	{
+	  for (k = 13; k < m3; k++)
+	    {
+	      if ((i + 2) * 12 + (j + 5) * 4 + (k - 13) !=  9 + f++)
+		l2++;
+	    }
+	}
+
     for (i = -2; i < m1; i++)
       for (j = m2; j < -2; j++)
 	{
@@ -30,7 +43,7 @@ main (void)
 	    }
 	}
 
-  if (l != x)
+  if (l != x || l2 != x)
     abort ();
 
   return 0;
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/collapse-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/collapse-3.c
index 680042892e4..50f538d0a32 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/collapse-3.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/collapse-3.c
@@ -7,7 +7,7 @@
 int
 main (void)
 {
-  int i2, l = 0, r = 0;
+  int i2, l = 0, r = 0, l2 = 0;
   int a[3][3][3];
 
   memset (a, '\0', sizeof (a));
@@ -27,13 +27,24 @@ main (void)
 		l += 1;
     }
 
+  /*  Test loop with >= condition.  */
+#pragma acc parallel
+    {
+      #pragma acc loop collapse(2) reduction(|:l2)
+	for (i2 = 0; i2 < 2; i2++)
+	  for (int j = 1; j >= 0; j--)
+	    for (int k = 0; k < 2; k++)
+	      if (a[i2][j][k] != i2 + j * 4 + k * 16)
+		l2 += 1;
+    }
+
     for (i2 = 0; i2 < 2; i2++)
       for (int j = 0; j < 2; j++)
 	for (int k = 0; k < 2; k++)
 	  if (a[i2][j][k] != i2 + j * 4 + k * 16)
 	    r += 1;
 
-  if (l != r)
+  if (l != r || l2 != r)
     abort ();
   return 0;
 }