[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