From patchwork Wed Oct 19 10:28:39 2016 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Thomas Schwinge X-Patchwork-Id: 684035 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Received: from sourceware.org (server1.sourceware.org [209.132.180.131]) (using TLSv1.2 with cipher ECDHE-RSA-AES256-GCM-SHA384 (256/256 bits)) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 3szSrM2zMfz9sCZ for ; Wed, 19 Oct 2016 21:30:26 +1100 (AEDT) Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b=b2oTvPtY; dkim-atps=neutral DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:from :to:cc:subject:in-reply-to:references:date:message-id :mime-version:content-type:content-transfer-encoding; q=dns; s= default; b=enZUMdEergRmAI3/EQamGCKbCIxD/Y/PTGhUBtrz+nKEgyLvoqmXh We0sI81rQV3TwvrJTTPTnvT1EcRzFyfSBoX7MgqHRrIG/RHUs9PQd+uHcWOE68IJ bKfqjE8Bfc5dtjX2Y87NX03R01Z9IOV/E5v2S155VKvd2MemZeUpZw= DKIM-Signature: v=1; a=rsa-sha1; c=relaxed; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:from :to:cc:subject:in-reply-to:references:date:message-id :mime-version:content-type:content-transfer-encoding; s=default; bh=/jdfL0resEfgOTwC391M4bWcoVY=; b=b2oTvPtYLKaVH1o170Z1lszyEJLv wjq7SU6zTL5xlDdKPgjy8nBCL2OBWvjiOuTbfJq65Wu0kVF25HFwmS2XxIL2aLrk fea8Y3yRDhGR21FVdhHNTFwjoGousIKNE0Uc5XnbtHrj4cRDyVhFl1xA5M5Q/jr5 k08Yvj+Pv24buj8= Received: (qmail 76960 invoked by alias); 19 Oct 2016 10:29:19 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Unsubscribe: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Delivered-To: mailing list gcc-patches@gcc.gnu.org Received: (qmail 76941 invoked by uid 89); 19 Oct 2016 10:29:18 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-1.9 required=5.0 tests=AWL, BAYES_00, RCVD_IN_DNSWL_NONE, SPF_PASS, URIBL_RED autolearn=ham version=3.3.2 spammy= X-Spam-User: qpsmtpd, 2 recipients X-HELO: relay1.mentorg.com Received: from relay1.mentorg.com (HELO relay1.mentorg.com) (192.94.38.131) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Wed, 19 Oct 2016 10:29:08 +0000 Received: from nat-ies.mentorg.com ([192.94.31.2] helo=svr-ies-mbx-01.mgc.mentorg.com) by relay1.mentorg.com with esmtp id 1bwo7N-0002zJ-Dz from Thomas_Schwinge@mentor.com ; Wed, 19 Oct 2016 03:29:05 -0700 Received: from hertz.schwinge.homeip.net (137.202.0.87) by svr-ies-mbx-01.mgc.mentorg.com (139.181.222.1) with Microsoft SMTP Server (TLS) id 15.0.1210.3; Wed, 19 Oct 2016 11:29:02 +0100 From: Thomas Schwinge To: GCC Patches , Richard Biener CC: Nathan Sidwell , GCC Development , Subject: [PR tree-optimization/78024] Clear basic block flags before using BB_VISITED for OpenACC loops processing In-Reply-To: References: <8760ov2wbg.fsf@hertz.schwinge.homeip.net> <9c7622ab-0c76-b054-ea4b-34df71febe35@acm.org> <87zim31fjx.fsf@hertz.schwinge.homeip.net> <87wph719jm.fsf@hertz.schwinge.homeip.net> <87twcb14et.fsf@hertz.schwinge.homeip.net> <87d1ixwi2v.fsf@hertz.schwinge.homeip.net> User-Agent: Notmuch/0.9-101-g81dad07 (http://notmuchmail.org) Emacs/24.4.1 (x86_64-pc-linux-gnu) Date: Wed, 19 Oct 2016 12:28:39 +0200 Message-ID: <8760oows2w.fsf@hertz.schwinge.homeip.net> MIME-Version: 1.0 X-ClientProxiedBy: svr-ies-mbx-01.mgc.mentorg.com (139.181.222.1) To svr-ies-mbx-01.mgc.mentorg.com (139.181.222.1) Hi! On Wed, 19 Oct 2016 12:07:13 +0200, Richard Biener wrote: > On Tue, Oct 18, 2016 at 9:52 PM, Thomas Schwinge > wrote: > > can I at > > least commit the OpenACC loops processing fix? Here is the latest > > version, simplified after your r241296 IRA vs. BB_VISITED fixes: > > Sure, I considered that approved already I don't see such an approval anywhere? > (it's even obvious). I've been told to be very careful in GCC with declaring something "obvious". (Though, I will be happy to increase my autonomy in declaring patches "obvious"/"ready for commit". This will potentially save me from wasting hours when keeping patches up-to-date, while repeatedly begging for review, and so on.) Anyway, without changes now committed in r241334: commit b3e3b38b4b1ac65df70d98f7e7557a27947948c1 Author: tschwinge Date: Wed Oct 19 10:19:24 2016 +0000 [PR tree-optimization/78024] Clear basic block flags before using BB_VISITED for OpenACC loops processing gcc/ * omp-low.c (oacc_loop_discovery): Call clear_bb_flags before, and don't clear BB_VISITED after processing. gcc/testsuite/ * gcc.dg/goacc/loop-processing-1.c: New file. git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@241334 138bc75d-0d04-0410-961f-82ee72b054a4 --- gcc/ChangeLog | 6 ++++++ gcc/omp-low.c | 8 +++----- gcc/testsuite/ChangeLog | 5 +++++ gcc/testsuite/gcc.dg/goacc/loop-processing-1.c | 18 ++++++++++++++++++ 4 files changed, 32 insertions(+), 5 deletions(-) Grüße Thomas diff --git gcc/ChangeLog gcc/ChangeLog index d5830d5..4117eb3 100644 --- gcc/ChangeLog +++ gcc/ChangeLog @@ -1,3 +1,9 @@ +2016-10-19 Thomas Schwinge + + PR tree-optimization/78024 + * omp-low.c (oacc_loop_discovery): Call clear_bb_flags before, and + don't clear BB_VISITED after processing. + 2016-10-19 Richard Biener * domwalk.c (dom_walker::walk): Use RPO order. diff --git gcc/omp-low.c gcc/omp-low.c index 77f89d5..3ef796f 100644 --- gcc/omp-low.c +++ gcc/omp-low.c @@ -19236,7 +19236,9 @@ oacc_loop_sibling_nreverse (oacc_loop *loop) static oacc_loop * oacc_loop_discovery () { - basic_block bb; + /* Clear basic block flags, in particular BB_VISITED which we're going to use + in the following. */ + clear_bb_flags (); oacc_loop *top = new_oacc_loop_outer (current_function_decl); oacc_loop_discover_walk (top, ENTRY_BLOCK_PTR_FOR_FN (cfun)); @@ -19245,10 +19247,6 @@ oacc_loop_discovery () that diagnostics come out in an unsurprising order. */ top = oacc_loop_sibling_nreverse (top); - /* Reset the visited flags. */ - FOR_ALL_BB_FN (bb, cfun) - bb->flags &= ~BB_VISITED; - return top; } diff --git gcc/testsuite/ChangeLog gcc/testsuite/ChangeLog index df73022..ec875aa 100644 --- gcc/testsuite/ChangeLog +++ gcc/testsuite/ChangeLog @@ -1,3 +1,8 @@ +2016-10-19 Thomas Schwinge + + PR tree-optimization/78024 + * gcc.dg/goacc/loop-processing-1.c: New file. + 2016-10-19 Richard Biener * gcc.dg/tree-ssa/pr61839_2.c: Fix testcase. diff --git gcc/testsuite/gcc.dg/goacc/loop-processing-1.c gcc/testsuite/gcc.dg/goacc/loop-processing-1.c new file mode 100644 index 0000000..619576a --- /dev/null +++ gcc/testsuite/gcc.dg/goacc/loop-processing-1.c @@ -0,0 +1,18 @@ +/* Make sure that OpenACC loop processing happens. */ +/* { dg-additional-options "-O2 -fdump-tree-oaccdevlow" } */ + +extern int place (); + +void vector_1 (int *ary, int size) +{ +#pragma acc parallel num_workers (32) vector_length(32) copy(ary[0:size]) firstprivate (size) + { +#pragma acc loop gang + for (int jx = 0; jx < 1; jx++) +#pragma acc loop auto + for (int ix = 0; ix < size; ix++) + ary[ix] = place (); + } +} + +/* { dg-final { scan-tree-dump {OpenACC loops.*Loop 0\(0\).*Loop 14\(1\).*\.data_dep\.[0-9_]+ = UNIQUE \(OACC_HEAD_MARK, 0, 1, 20\);.*Head-0:.*\.data_dep\.[0-9_]+ = UNIQUE \(OACC_HEAD_MARK, 0, 1, 20\);.*\.data_dep\.[0-9_]+ = UNIQUE \(OACC_FORK, \.data_dep\.[0-9_]+, 0\);.*Tail-0:.*\.data_dep\.[0-9_]+ = UNIQUE \(OACC_TAIL_MARK, \.data_dep\.[0-9_]+, 1\);.*\.data_dep\.[0-9_]+ = UNIQUE \(OACC_JOIN, \.data_dep\.[0-9_]+, 0\);.*Loop 6\(4\).*\.data_dep\.[0-9_]+ = UNIQUE \(OACC_HEAD_MARK, 0, 1, 6\);.*Head-0:.*\.data_dep\.[0-9_]+ = UNIQUE \(OACC_HEAD_MARK, 0, 1, 6\);.*\.data_dep\.[0-9_]+ = UNIQUE \(OACC_FORK, \.data_dep\.[0-9_]+, 2\);.*Tail-0:.*\.data_dep\.[0-9_]+ = UNIQUE \(OACC_TAIL_MARK, \.data_dep\.[0-9_]+, 1\);.*\.data_dep\.[0-9_]+ = UNIQUE \(OACC_JOIN, \.data_dep\.[0-9_]+, 2\);} "oaccdevlow" } } */