public inbox for gcc-cvs@sourceware.org help / color / mirror / Atom feed
From: Hafiz Abid Qadeer <abidh@gcc.gnu.org> To: gcc-cvs@gcc.gnu.org Subject: [gcc/devel/omp/gcc-10] [OpenACC] Fix an ICE where a loop with GT condition is collapsed. Date: Sun, 11 Apr 2021 15:39:21 +0000 (GMT) [thread overview] Message-ID: <20210411153921.02DD73844041@sourceware.org> (raw) 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; }
reply other threads:[~2021-04-11 15:39 UTC|newest] Thread overview: [no followups] expand[flat|nested] mbox.gz Atom feed
Reply instructions: You may reply publicly to this message via plain-text email using any one of the following methods: * Save the following mbox file, import it into your mail client, and reply-to-all from there: mbox Avoid top-posting and favor interleaved quoting: https://en.wikipedia.org/wiki/Posting_style#Interleaved_style * Reply using the --to, --cc, and --in-reply-to switches of git-send-email(1): git send-email \ --in-reply-to=20210411153921.02DD73844041@sourceware.org \ --to=abidh@gcc.gnu.org \ --cc=gcc-cvs@gcc.gnu.org \ /path/to/YOUR_REPLY https://kernel.org/pub/software/scm/git/docs/git-send-email.html * If your mail client supports setting the In-Reply-To header via mailto: links, try the mailto: linkBe sure your reply has a Subject: header at the top and a blank line before the message body.
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).