From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: (qmail 49692 invoked by alias); 30 Jan 2020 21:13:19 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Received: (qmail 49642 invoked by uid 89); 30 Jan 2020 21:13:19 -0000 Authentication-Results: sourceware.org; auth=none X-Spam-SWARE-Status: No, score=-18.9 required=5.0 tests=AWL,BAYES_00,GIT_PATCH_0,GIT_PATCH_1,GIT_PATCH_2,GIT_PATCH_3,KAM_SHORT,SPF_PASS autolearn=ham version=3.3.1 spammy=serial, PR93488, user-visible, valid-code X-HELO: esa2.mentor.iphmx.com Received: from esa2.mentor.iphmx.com (HELO esa2.mentor.iphmx.com) (68.232.141.98) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Thu, 30 Jan 2020 21:13:17 +0000 IronPort-SDR: QP6vyzWcIvwwpBTUb2jXOljM/5VTBn0YIEXaKTjhfHvC/gOn6hXIu6plcDM9rH/c40uMbcHqnz djp5DHlhWhIj3dQel+L7xh5BoXxilj0mnE9Sb1FrscIOCsf2LsofXvANoUtxEsApC8Wi6NdaHz dkAZhdWcH2nIcgDRchmapkELpSYWscP7/K/I7eB3e8FcP9jZAWXwc5+cHVBc+ZAsKzVy6m5QAa 31fIXTkrUoJphdrSO+71vAConvFLIPKhsIJo1BTuLro1NMcZ/ieyyewMo8AsxNE8XD7B42huj1 b18= Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa2.mentor.iphmx.com with ESMTP; 30 Jan 2020 13:13:16 -0800 IronPort-SDR: s+0KnGQ1DZjLKG3UnKHZgZx7bBRpV+YHnVi8R47DUw65T/lmnZy21ezum9lUGJcTmTGxL/CF8W ksF6KB/KY+WA== Subject: Re: [PR93488] [OpenACC] ICE in type-cast 'async', 'wait' clauses To: Thomas Schwinge CC: , Jakub Jelinek References: <87r1zimomb.fsf@euler.schwinge.homeip.net> From: Andrew Stubbs Message-ID: Date: Fri, 31 Jan 2020 00:05:00 -0000 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:68.0) Gecko/20100101 Thunderbird/68.4.1 MIME-Version: 1.0 In-Reply-To: <87r1zimomb.fsf@euler.schwinge.homeip.net> Content-Type: multipart/mixed; boundary="------------9325DF3CAEAE190F38C3094D" Return-Path: ams@codesourcery.com X-SW-Source: 2020-01/txt/msg02041.txt.bz2 --------------9325DF3CAEAE190F38C3094D Content-Type: text/plain; charset="utf-8"; format=flowed Content-Transfer-Encoding: 7bit Content-length: 1042 On 29/01/2020 12:30, Thomas Schwinge wrote: > As can be seen in the code a few lines below, the very same problem also > exists for the 'wait' clause; it seems reasonable to fix both at the same > time. This is not a recent regression, but a user-visible valid-code ICE > that has existed "forever"; I filed "ICE in > type-cast 'async', 'wait' clauses" for tracking. This problem is similar > to the OpenMP 'device' clause "ICE in > verify_gimple_in_cfg, at tree-cfg.c:5212"; I suggest we also use > 'force_gimple_operand_gsi' instead of manually doing your suggested > 'create_tmp_var', 'gimple_build_assign', 'gsi_insert_before'. Include a > test case that covers all relevant code paths; I've attached a test case > to the PR, but I've not verified whether "that covers *all* relevant code > paths". This should then be backported to all GCC release branches; I > can easily test the backports for you, if you're not already set up to do > such testing. How's this? Andrew --------------9325DF3CAEAE190F38C3094D Content-Type: text/x-patch; charset="UTF-8"; name="200130-normalize-goacc-args.patch" Content-Transfer-Encoding: 7bit Content-Disposition: inline; filename="200130-normalize-goacc-args.patch" Content-length: 2152 Normalize GOACC_parallel_keyed async and wait parameters 2020-01-30 Andrew Stubbs Thomas Schwinge PR middle-end/93488 gcc/ * omp-expand.c (expand_omp_target): Use force_gimple_operand_gsi on t_async and the wait arguments. diff --git a/gcc/omp-expand.c b/gcc/omp-expand.c index cd423ad799e..ec4baf46965 100644 --- a/gcc/omp-expand.c +++ b/gcc/omp-expand.c @@ -8418,7 +8418,9 @@ expand_omp_target (struct omp_region *region) i_async)); } if (t_async) - args.safe_push (t_async); + args.safe_push (force_gimple_operand_gsi (&gsi, t_async, true, + NULL_TREE, true, + GSI_SAME_STMT)); /* Save the argument index, and ... */ unsigned t_wait_idx = args.length (); @@ -8431,9 +8433,12 @@ expand_omp_target (struct omp_region *region) for (; c; c = OMP_CLAUSE_CHAIN (c)) if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_WAIT) { - args.safe_push (fold_convert_loc (OMP_CLAUSE_LOCATION (c), - integer_type_node, - OMP_CLAUSE_WAIT_EXPR (c))); + tree arg = fold_convert_loc (OMP_CLAUSE_LOCATION (c), + integer_type_node, + OMP_CLAUSE_WAIT_EXPR (c)); + arg = force_gimple_operand_gsi (&gsi, arg, true, NULL_TREE, true, + GSI_SAME_STMT); + args.safe_push (arg); num_waits++; } diff --git a/gcc/testsuite/c-c++-common/goacc/pr93488.c b/gcc/testsuite/c-c++-common/goacc/pr93488.c new file mode 100644 index 00000000000..6fddad919d2 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/pr93488.c @@ -0,0 +1,22 @@ +/* PR middle-end/93488 + + Ensure that wait and async arguments can be cast to the correct type + without breaking gimple verification. */ + +void test() +{ + /* int */ unsigned char a = 1; + /* int */ unsigned char w = 1; + +#pragma acc parallel wait(w) async(a) + ; +#pragma acc kernels wait(w) async(a) + ; +#pragma acc serial wait(w) async(a) + ; + int data = 0; +#pragma acc enter data wait(w) async(a) create(data) +#pragma acc update wait(w) async(a) device(data) +#pragma acc exit data wait(w) async(a) delete(data) +#pragma acc wait(w) async(a) +} --------------9325DF3CAEAE190F38C3094D--