[PATCH] [OpenACC] Fix an ICE where a loop with GT condition is collapsed.

[PATCH] [OpenACC] Fix an ICE where a loop with GT condition is collapsed.

[Hafiz Abid Qadeer]

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;
#pragma acc parallel
  {
#pragma acc loop independent gang, collapse(2)
       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.

	gcc/
	* omp-expand.c (expand_oacc_collapse_init): Update condition in
	a gcc_assert.
	* 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.

---
 gcc/omp-expand.c                              |  2 +-
 gcc/testsuite/c-c++-common/goacc/collapse-2.c | 34 +++++++++++++++++++
 .../libgomp.oacc-c-c++-common/collapse-2.c    | 17 ++++++++--
 .../libgomp.oacc-c-c++-common/collapse-3.c    | 15 ++++++--
 4 files changed, 63 insertions(+), 5 deletions(-)
 create mode 100644 gcc/testsuite/c-c++-common/goacc/collapse-2.c

diff --git a/gcc/omp-expand.c b/gcc/omp-expand.c
index 7559ec80263..dc797f95154 100644
--- a/gcc/omp-expand.c
+++ b/gcc/omp-expand.c
@@ -1541,7 +1541,7 @@ expand_oacc_collapse_init (const struct omp_for_data *fd,
       tree iter_type = TREE_TYPE (loop->v);
       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..97328960932
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/collapse-2.c
@@ -0,0 +1,34 @@
+/* Test for ICE when loop with > condition is being collapsed.  */
+/* { dg-skip-if "not yet" { c++ } } */
+
+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--)
+	  ;
+}
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;
 }
-- 
2.25.1

Re: [PATCH] [OpenACC] Fix an ICE where a loop with GT condition is collapsed.

[Thomas Schwinge]

Hi Abid!

Is <abidh@sourceware.org> enabled for accessing GCC (that is,
<abidh@gcc.gnu.org>)?  If not, please request as indicated on
<https://sourceware.org/cgi-bin/pdw/ps_form.cgi>: "send an email to the
overseers mail account at this site telling what project you want write
access to and who approved that access".


On 2021-04-09T11:38:57+0100, Hafiz Abid Qadeer <abidh@codesourcery.com> wrote:
> 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;
> #pragma acc parallel
>   {
> #pragma acc loop independent gang, collapse(2)
>        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.

Thanks for the description with sufficient analysis and detail!  I'm not
very familiar with the 'collapse' handling, but yes, this makes sense.

Given that it's the same 'gcc/omp-expand.c:expand_oacc_collapse_init'
code and 'assert', I suppose this is (and your patch fixes) the very same
issues as discussed in <https://gcc.gnu.org/PR98088> "ICE in
expand_oacc_collapse_init, at omp-expand.c:1533".  Will you please verify
that, and if yes, assign that one to you (<abidh@gcc.gnu.org>), add
PR98088 marker to your Git commit log and ChangeLog update therein (see
existing ones for reference), and also include the testcases as discussed
in PR98088 (even the "invalid" ones, to make sure the ICE goes away).

> I have fixed it by modifying the check in the assertion to be same as
> check on fd->loop->cond_code.

That seems to match Jakub's 2021-02-01 comment PR98088, good.

> 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.
>
>       gcc/
>       * omp-expand.c (expand_oacc_collapse_init): Update condition in
>       a gcc_assert.
>       * 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.

ACK, thanks.  This then qualifies for all GCC release branches.  (I can
easily handle the backports for you, if you'd like.)


Grüße
 Thomas


>  gcc/omp-expand.c                              |  2 +-
>  gcc/testsuite/c-c++-common/goacc/collapse-2.c | 34 +++++++++++++++++++
>  .../libgomp.oacc-c-c++-common/collapse-2.c    | 17 ++++++++--
>  .../libgomp.oacc-c-c++-common/collapse-3.c    | 15 ++++++--
>  4 files changed, 63 insertions(+), 5 deletions(-)
>  create mode 100644 gcc/testsuite/c-c++-common/goacc/collapse-2.c
>
> diff --git a/gcc/omp-expand.c b/gcc/omp-expand.c
> index 7559ec80263..dc797f95154 100644
> --- a/gcc/omp-expand.c
> +++ b/gcc/omp-expand.c
> @@ -1541,7 +1541,7 @@ expand_oacc_collapse_init (const struct omp_for_data *fd,
>        tree iter_type = TREE_TYPE (loop->v);
>        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..97328960932
> --- /dev/null
> +++ b/gcc/testsuite/c-c++-common/goacc/collapse-2.c
> @@ -0,0 +1,34 @@
> +/* Test for ICE when loop with > condition is being collapsed.  */
> +/* { dg-skip-if "not yet" { c++ } } */
> +
> +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--)
> +       ;
> +}
> 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;
>  }
> --
> 2.25.1
-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstrasse 201, 80634 München Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Frank Thürauf

Re: [PATCH] [OpenACC] Fix an ICE where a loop with GT condition is collapsed.

[Tobias Burnus]

Hi Abid,

I think that's the same issue as https://gcc.gnu.org/PR98088
if so, please add 'PR middle-end/98088' to the changelog.

I think the second testcase is covered, but you could also add
the first testcase to c-c++-common/goacc/collapse-2.c – the first
testcase in the PR is for a zero-trip loop.

(more notes below)

On 09.04.21 12:38, Hafiz Abid Qadeer wrote:
> 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;
> #pragma acc parallel
>    {
> #pragma acc loop independent gang, collapse(2)
>         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.
>
>       gcc/
>       * omp-expand.c (expand_oacc_collapse_init): Update condition in
>       a gcc_assert.
>       * testsuite/c-c++-common/goacc/collapse-2.c: New.

The changelog file is in 'gcc/testsuite/', hence, you need a separate
entry (and remove 'testsuite/').

>       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.
>
> ---
>   gcc/omp-expand.c                              |  2 +-
>   gcc/testsuite/c-c++-common/goacc/collapse-2.c | 34 +++++++++++++++++++
>   .../libgomp.oacc-c-c++-common/collapse-2.c    | 17 ++++++++--
>   .../libgomp.oacc-c-c++-common/collapse-3.c    | 15 ++++++--
>   4 files changed, 63 insertions(+), 5 deletions(-)
>   create mode 100644 gcc/testsuite/c-c++-common/goacc/collapse-2.c
>
> diff --git a/gcc/omp-expand.c b/gcc/omp-expand.c
> index 7559ec80263..dc797f95154 100644
> --- a/gcc/omp-expand.c
> +++ b/gcc/omp-expand.c
> @@ -1541,7 +1541,7 @@ expand_oacc_collapse_init (const struct omp_for_data *fd,
>         tree iter_type = TREE_TYPE (loop->v);
>         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..97328960932
> --- /dev/null
> +++ b/gcc/testsuite/c-c++-common/goacc/collapse-2.c
> @@ -0,0 +1,34 @@
> +/* Test for ICE when loop with > condition is being collapsed.  */
> +/* { dg-skip-if "not yet" { c++ } } */

Why is the dg-skip-if needed? Codewise it does not seem to be required
and it also does compile with g++ after applying the patch.

With this removed and the other nits fixed: LGTM & thanks for the patch.

After some grace time, please also apply to GCC 10 (and possibly GCC 9)
besides OG10, given that it is a regression and a trivial fix.

Tobias

> +
> +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--)
> +       ;
> +}
> 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;
>   }
-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstrasse 201, 80634 München Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Frank Thürauf

Re: [PATCH] [OpenACC] Fix an ICE where a loop with GT condition is collapsed.

[Hafiz Abid Qadeer]

[-- Attachment #1: Type: text/plain, Size: 1066 bytes --]

Thanks for the review. Please see my comments below.

On 09/04/2021 13:48, Tobias Burnus wrote:
> Hi Abid,
> 
> I think that's the same issue as https://gcc.gnu.org/PR98088
> if so, please add 'PR middle-end/98088' to the changelog.
Done
> 
> I think the second testcase is covered, but you could also add
> the first testcase to c-c++-common/goacc/collapse-2.c – the first
> testcase in the PR is for a zero-trip loop.
I have added the tests from the PR.

>>     * testsuite/c-c++-common/goacc/collapse-2.c: New.
> 
> The changelog file is in 'gcc/testsuite/', hence, you need a separate entry (and
> remove 'testsuite/').
Done

>> +/* { dg-skip-if "not yet" { c++ } } */
> 
> Why is the dg-skip-if needed? Codewise it does not seem to be required
> and it also does compile with g++ after applying the patch.
It was in the collapse-1.c file. I have removed it now.

> 
> With this removed and the other nits fixed: LGTM & thanks for the patch.

Updated patch attached. I will apply if there are no further comments and once I
get write access.

Thanks,
Abid

[-- Attachment #2: 0001-OpenACC-Fix-an-ICE-where-a-loop-with-GT-condition-is.patch --]
[-- Type: text/x-patch, Size: 6117 bytes --]

From ecd9e83cce37d78b17dc186a9b75f2ff9e35eeb0 Mon Sep 17 00:00:00 2001
From: Hafiz Abid Qadeer <abidh@codesourcery.com>
Date: Thu, 8 Apr 2021 17:31:30 +0100
Subject: [PATCH] [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.
---
 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(-)
 create mode 100644 gcc/testsuite/c-c++-common/goacc/collapse-2.c

diff --git a/gcc/omp-expand.c b/gcc/omp-expand.c
index 7559ec80263..dc797f95154 100644
--- a/gcc/omp-expand.c
+++ b/gcc/omp-expand.c
@@ -1541,7 +1541,7 @@ expand_oacc_collapse_init (const struct omp_for_data *fd,
       tree iter_type = TREE_TYPE (loop->v);
       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;
 }
-- 
2.25.1

Re: [PATCH] [OpenACC] Fix an ICE where a loop with GT condition is collapsed.

[Christophe Lyon]

Hi,



On Fri, 9 Apr 2021 at 18:44, Hafiz Abid Qadeer <abid_qadeer@mentor.com> wrote:
>
> Thanks for the review. Please see my comments below.
>
> On 09/04/2021 13:48, Tobias Burnus wrote:
> > Hi Abid,
> >
> > I think that's the same issue as https://gcc.gnu.org/PR98088
> > if so, please add 'PR middle-end/98088' to the changelog.
> Done
> >
> > I think the second testcase is covered, but you could also add
> > the first testcase to c-c++-common/goacc/collapse-2.c – the first
> > testcase in the PR is for a zero-trip loop.
> I have added the tests from the PR.
>
> >>     * testsuite/c-c++-common/goacc/collapse-2.c: New.
> >
> > The changelog file is in 'gcc/testsuite/', hence, you need a separate entry (and
> > remove 'testsuite/').
> Done
>
> >> +/* { dg-skip-if "not yet" { c++ } } */
> >
> > Why is the dg-skip-if needed? Codewise it does not seem to be required
> > and it also does compile with g++ after applying the patch.
> It was in the collapse-1.c file. I have removed it now.
>
> >
> > With this removed and the other nits fixed: LGTM & thanks for the patch.
>
> Updated patch attached. I will apply if there are no further comments and once I
> get write access.
>

The new test generates an ICE on aarch64-linux-gnu in the gcc-10 branch:
/gcc/testsuite/c-c++-common/goacc/collapse-2.c: In function 'f5._omp_fn.0':
/gcc/testsuite/c-c++-common/goacc/collapse-2.c:56:1: error: type
mismatch in binary expression
signed long

long int

int

D.3972 = .tile.57 * .tile.59;
during IPA pass: *free_lang_data
/gcc/testsuite/c-c++-common/goacc/collapse-2.c:56:1: internal compiler
error: verify_gimple failed
0xdc554e verify_gimple_in_cfg(function*, bool)
        /gcc/tree-cfg.c:5502
0xc7e7f2 execute_function_todo
        /gcc/passes.c:1985
0xc7f0bd do_per_function
        /gcc/passes.c:1647
0xc7f195 execute_todo
        /gcc/passes.c:2039
Please submit a full bug report,
with preprocessed source if appropriate.
Please include the complete backtrace with any bug report.
See <https://gcc.gnu.org/bugs/> for instructions.
compiler exited with status 1
FAIL: c-c++-common/goacc/collapse-2.c (internal compiler error)
FAIL: c-c++-common/goacc/collapse-2.c 3 blank line(s) in output
FAIL: c-c++-common/goacc/collapse-2.c (test for excess errors)



Can you check?

Thanks

Christophe

> Thanks,
> Abid

Re: [PATCH] [OpenACC] Fix an ICE where a loop with GT condition is collapsed.

[Tobias Burnus]

On 19.04.21 10:48, Christophe Lyon wrote:

> The new test generates an ICE on aarch64-linux-gnu in the gcc-10 branch:
Looks as someone (like me) should backport https://gcc.gnu.org/97880 /
r11-5489 to GCC 10.
> /gcc/testsuite/c-c++-common/goacc/collapse-2.c: In function 'f5._omp_fn.0':
> /gcc/testsuite/c-c++-common/goacc/collapse-2.c:56:1: error: type
> mismatch in binary expression
> signed long
>
> long int
>
> int
>
> D.3972 = .tile.57 * .tile.59;
> during IPA pass: *free_lang_data

Tobias

-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstrasse 201, 80634 München Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Frank Thürauf

Re: [PATCH] [OpenACC] Fix an ICE where a loop with GT condition is collapsed.

[Tobias Burnus]

On 19.04.21 11:25, Tobias Burnus wrote:
> On 19.04.21 10:48, Christophe Lyon wrote:
>> The new test generates an ICE on aarch64-linux-gnu in the gcc-10 branch:
> Looks as someone (like me) should backport https://gcc.gnu.org/97880 /
> r11-5489 to GCC 10.
("OpenACC: Fix integer-type issue with collapse/tile [PR97880]")

Now backported to GCC 10 as
r10-9716-gaf408874e3d94492ac08ba17462b3ff8ecb4d791, please confirm that
it did fix your issue:

>> /gcc/testsuite/c-c++-common/goacc/collapse-2.c: In function
>> 'f5._omp_fn.0':
>> /gcc/testsuite/c-c++-common/goacc/collapse-2.c:56:1: error: type
>> mismatch in binary expression
>> signed long
>>
>> long int
>>
>> int
>>
>> D.3972 = .tile.57 * .tile.59;
>> during IPA pass: *free_lang_data
Tobias
-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstrasse 201, 80634 München Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Frank Thürauf

Re: [PATCH] [OpenACC] Fix an ICE where a loop with GT condition is collapsed.

[Christophe Lyon]

On Mon, 19 Apr 2021 at 12:40, Tobias Burnus <tobias@codesourcery.com> wrote:
>
> On 19.04.21 11:25, Tobias Burnus wrote:
> > On 19.04.21 10:48, Christophe Lyon wrote:
> >> The new test generates an ICE on aarch64-linux-gnu in the gcc-10 branch:
> > Looks as someone (like me) should backport https://gcc.gnu.org/97880 /
> > r11-5489 to GCC 10.
> ("OpenACC: Fix integer-type issue with collapse/tile [PR97880]")
>
> Now backported to GCC 10 as
> r10-9716-gaf408874e3d94492ac08ba17462b3ff8ecb4d791, please confirm that
> it did fix your issue:
>

Yes, that works, thanks

Christophe

> >> /gcc/testsuite/c-c++-common/goacc/collapse-2.c: In function
> >> 'f5._omp_fn.0':
> >> /gcc/testsuite/c-c++-common/goacc/collapse-2.c:56:1: error: type
> >> mismatch in binary expression
> >> signed long
> >>
> >> long int
> >>
> >> int
> >>
> >> D.3972 = .tile.57 * .tile.59;
> >> during IPA pass: *free_lang_data
> Tobias
> -----------------
> Mentor Graphics (Deutschland) GmbH, Arnulfstrasse 201, 80634 München Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Frank Thürauf