From patchwork Fri Apr 9 10:38:57 2021 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Abid Qadeer X-Patchwork-Id: 1464275 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (sender SPF authorized) smtp.mailfrom=gcc.gnu.org (client-ip=8.43.85.97; helo=sourceware.org; envelope-from=gcc-patches-bounces@gcc.gnu.org; receiver=) Received: from sourceware.org (server2.sourceware.org [8.43.85.97]) (using TLSv1.3 with cipher TLS_AES_256_GCM_SHA384 (256/256 bits) key-exchange X25519 server-signature RSA-PSS (4096 bits) server-digest SHA256) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 4FGvkt43Qfz9sSC for ; Fri, 9 Apr 2021 20:39:16 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id AF639385DC2E; Fri, 9 Apr 2021 10:39:12 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa3.mentor.iphmx.com (esa3.mentor.iphmx.com [68.232.137.180]) by sourceware.org (Postfix) with ESMTPS id DAFCD385E017 for ; Fri, 9 Apr 2021 10:39:08 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.3.2 sourceware.org DAFCD385E017 Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=abid_qadeer@mentor.com IronPort-SDR: lW90k88jXbmIsfF2oC1yNiGJzrGUn8uzQHUVtLqec7krTVYtLwqP91NtSi+lGpc3nKl5+FFX7h 6qtVnt2CJRTRvEO1KfdebyPphtnAX28Qm5ePwzhpvKjd5ns4oFD6z5SPbhzLc9f4DJepPXCm8O svurNxgPzDIwfTkAZEG5UNr9fWO8MvYejevD0KR5Mq9NrkTXB3r47qG4D0fNTxHqJCrqFabdfR /3DuQ90LknzPTKWJZkjayi99CRBXYMrV122Yc6B8PzDi/MHJXYM3hVGbXayDITz25PHTIOISFf C6Y= X-IronPort-AV: E=Sophos;i="5.82,209,1613462400"; d="scan'208";a="59977722" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa3.mentor.iphmx.com with ESMTP; 09 Apr 2021 02:39:07 -0800 IronPort-SDR: F5q/KFaCkAs9ZHoBjTkEq+HS5L0kChs/T13kVbOEOZq9bul4ytAL4anxFw3o4DJ7Ha7g8Wnmej Tmro1W/7eyJ8L9ZRZxPh/naito/7MLBPlZB47KR/MpPVACcB2wZYSyHlng+7VnwJ60Q7O+mB1C hG/o2aZpviPknsJv6W/RzpZons2/+O2eOkI+WIqfI+o32UGUAcFzJh1TKGu089/1GWiDW10TWu Fk5YS5elnFdgVvpOpGuPe9nWGkf9ORz5D+cIHxW7hPY87sPw5RofsB6mR6mNjwfrM+RZszv3VV V2E= From: Hafiz Abid Qadeer To: Subject: [PATCH] [OpenACC] Fix an ICE where a loop with GT condition is collapsed. Date: Fri, 9 Apr 2021 11:38:57 +0100 Message-ID: <20210409103857.2707663-1-abidh@codesourcery.com> X-Mailer: git-send-email 2.25.1 MIME-Version: 1.0 X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: SVR-IES-MBX-04.mgc.mentorg.com (139.181.222.4) To SVR-IES-MBX-03.mgc.mentorg.com (139.181.222.3) X-Spam-Status: No, score=-13.5 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, SPF_HELO_PASS, SPF_PASS, TXREP autolearn=ham autolearn_force=no version=3.4.2 X-Spam-Checker-Version: SpamAssassin 3.4.2 (2018-09-13) on server2.sourceware.org X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Cc: tobias@codesourcery.com, thomas@codesourcery.com, abidh@codesourcery.com Errors-To: gcc-patches-bounces@gcc.gnu.org Sender: "Gcc-patches" 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. 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. 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. --- 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; }