public inbox for gcc-cvs@sourceware.org
help / color / mirror / Atom feed
* [gcc/devel/omp/gcc-10] [OpenACC] Fix an ICE where a loop with GT condition is collapsed.
@ 2021-04-11 15:39 Hafiz Abid Qadeer
0 siblings, 0 replies; only message in thread
From: Hafiz Abid Qadeer @ 2021-04-11 15:39 UTC (permalink / raw)
To: gcc-cvs
https://gcc.gnu.org/g:bf869eaeaa56a9b56f035d4a67541a7cb85c6e0b
commit bf869eaeaa56a9b56f035d4a67541a7cb85c6e0b
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.
Diff:
---
gcc/ChangeLog.omp | 9 ++++
gcc/omp-expand.c | 2 +-
gcc/testsuite/ChangeLog.omp | 8 ++++
gcc/testsuite/c-c++-common/goacc/collapse-2.c | 56 ++++++++++++++++++++++
libgomp/ChangeLog.omp | 10 ++++
.../libgomp.oacc-c-c++-common/collapse-2.c | 17 ++++++-
.../libgomp.oacc-c-c++-common/collapse-3.c | 15 +++++-
7 files changed, 112 insertions(+), 5 deletions(-)
diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp
index d6566ce97a0..a89d17959c0 100644
--- a/gcc/ChangeLog.omp
+++ b/gcc/ChangeLog.omp
@@ -1,3 +1,12 @@
+2021-04-11 Hafiz Abid Qadeer <abidh@codesourcery.com>
+
+ Backport from mainline
+ 2021-04-11 Hafiz Abid Qadeer <abidh@codesourcery.com>
+
+ PR middle-end/98088
+ * omp-expand.c (expand_oacc_collapse_init): Update condition in
+ a gcc_assert.
+
2021-03-25 Kwok Cheung Yeung <kcy@codesourcery.com>
Backport from mainline
diff --git a/gcc/omp-expand.c b/gcc/omp-expand.c
index f8347c0dc59..c01351d26a9 100644
--- a/gcc/omp-expand.c
+++ b/gcc/omp-expand.c
@@ -1583,7 +1583,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/ChangeLog.omp b/gcc/testsuite/ChangeLog.omp
index 97cad18a069..45277a37fd2 100644
--- a/gcc/testsuite/ChangeLog.omp
+++ b/gcc/testsuite/ChangeLog.omp
@@ -1,3 +1,11 @@
+2021-04-11 Hafiz Abid Qadeer <abidh@codesourcery.com>
+
+ Backport from mainline
+ 2021-04-11 Hafiz Abid Qadeer <abidh@codesourcery.com>
+
+ PR middle-end/98088
+ * c-c++-common/goacc/collapse-2.c: New.
+
2021-03-25 Thomas Schwinge <thomas@codesourcery.com>
* gfortran.dg/gomp/order-4.f90: Adjust.
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/ChangeLog.omp b/libgomp/ChangeLog.omp
index f131c2c79b7..492d49c35a5 100644
--- a/libgomp/ChangeLog.omp
+++ b/libgomp/ChangeLog.omp
@@ -1,3 +1,13 @@
+2021-04-11 Hafiz Abid Qadeer <abidh@codesourcery.com>
+
+ Backport from mainline
+ 2021-04-11 Hafiz Abid Qadeer <abidh@codesourcery.com>
+
+ PR middle-end/98088
+ * 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.
+
2021-03-26 Thomas Schwinge <thomas@codesourcery.com>
* testsuite/libgomp.oacc-fortran/derivedtypes-arrays-1.f90: Adjust
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;
}
^ permalink raw reply [flat|nested] only message in thread
only message in thread, other threads:[~2021-04-11 15:39 UTC | newest]
Thread overview: (only message) (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2021-04-11 15:39 [gcc/devel/omp/gcc-10] [OpenACC] Fix an ICE where a loop with GT condition is collapsed Hafiz Abid Qadeer
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).