From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: by sourceware.org (Postfix, from userid 1820) id 86AF13854835; Sun, 18 Apr 2021 11:56:10 +0000 (GMT) DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org 86AF13854835 MIME-Version: 1.0 Content-Transfer-Encoding: 7bit Content-Type: text/plain; charset="utf-8" From: Hafiz Abid Qadeer To: gcc-cvs@gcc.gnu.org Subject: [gcc r10-9711] [OpenACC] Fix an ICE where a loop with GT condition is collapsed. X-Act-Checkin: gcc X-Git-Author: Hafiz Abid Qadeer X-Git-Refname: refs/heads/releases/gcc-10 X-Git-Oldrev: b18305e6746b57ac051386465f6077700932ffc4 X-Git-Newrev: e4dcb3383bff4c209a918127551cabc56b4171ae Message-Id: <20210418115610.86AF13854835@sourceware.org> Date: Sun, 18 Apr 2021 11:56:10 +0000 (GMT) X-BeenThere: gcc-cvs@gcc.gnu.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: Gcc-cvs mailing list List-Unsubscribe: , List-Archive: List-Help: List-Subscribe: , X-List-Received-Date: Sun, 18 Apr 2021 11:56:10 -0000 https://gcc.gnu.org/g:e4dcb3383bff4c209a918127551cabc56b4171ae commit r10-9711-ge4dcb3383bff4c209a918127551cabc56b4171ae Author: Hafiz Abid Qadeer 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. (cherry picked from commit ac200799acb5cd2fb9e1758f6bf5fff1978daaeb) Diff: --- 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(-) diff --git a/gcc/omp-expand.c b/gcc/omp-expand.c index 90df5c90547..b73fcf0b0a2 100644 --- a/gcc/omp-expand.c +++ b/gcc/omp-expand.c @@ -1569,7 +1569,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/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; }