ac200799ac
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.
51 lines
936 B
C
51 lines
936 B
C
/* { dg-do run } */
|
|
|
|
#include <stdlib.h>
|
|
|
|
int
|
|
main (void)
|
|
{
|
|
int i, j, k, l = 0, f = 0, x = 0, l2 = 0;
|
|
int m1 = 4, m2 = -5, m3 = 17;
|
|
|
|
#pragma acc parallel
|
|
#pragma acc loop seq collapse(3) reduction(+:l)
|
|
for (i = -2; i < m1; i++)
|
|
for (j = m2; j < -2; j++)
|
|
{
|
|
for (k = 13; k < m3; k++)
|
|
{
|
|
if ((i + 2) * 12 + (j + 5) * 4 + (k - 13) != 9 + f++)
|
|
l++;
|
|
}
|
|
}
|
|
|
|
/* 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++)
|
|
{
|
|
for (k = 13; k < m3; k++)
|
|
{
|
|
if ((i + 2) * 12 + (j + 5) * 4 + (k - 13) != 9 + f++)
|
|
x++;
|
|
}
|
|
}
|
|
|
|
if (l != x || l2 != x)
|
|
abort ();
|
|
|
|
return 0;
|
|
}
|