From patchwork Wed Aug 17 12:17:36 2016 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Chung-Lin Tang X-Patchwork-Id: 660090 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 3sDpCY1Cyfz9t22 for ; Wed, 17 Aug 2016 22:18:00 +1000 (AEST) Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b=qOwQLHIF; 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 :subject:to:references:cc:from:message-id:date:mime-version :in-reply-to:content-type; q=dns; s=default; b=WXTQplzdGY8Zvqtpx bwxeyyVmQescUP5713ewrA2VR011hG3TtyxMWVsjUuZB+b0OzIImjQJWnZ4dZ+42 TCSbCUsU9Bv660upEFhx8EO78BS2zOww9l4IXoLBQdrtFbXpLQM7JLpa5XZ6s8fY B3pmraJ2G7r/BWkhzk1WU9jxYQ= 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 :subject:to:references:cc:from:message-id:date:mime-version :in-reply-to:content-type; s=default; bh=llcDtgb8lv3iRgX2AvKbVXb 34+U=; b=qOwQLHIFTUh2AJY25BVdAdFlM28wCDIos0IKNpEFTGg/LK1dT/oqSi3 +m3LNYmS+ZFhvlPBIFCnHkCDWxM3NfRqc+r1DymKFduZsXU1hy32IV0JK9d1ZhyN zp4k3Qhnl6wPobC5uIgJA6kb1OVwecFDAbWZXhS4Ah/xlkK8WzJY= Received: (qmail 85781 invoked by alias); 17 Aug 2016 12:17:52 -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 85769 invoked by uid 89); 17 Aug 2016 12:17:52 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-1.1 required=5.0 tests=AWL, BAYES_00, KAM_ASCII_DIVIDERS, RCVD_IN_DNSWL_NONE, SPF_PASS autolearn=no version=3.3.2 spammy= 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, 17 Aug 2016 12:17:42 +0000 Received: from svr-orw-fem-03.mgc.mentorg.com ([147.34.97.39]) by relay1.mentorg.com with esmtp id 1bZzms-00064L-Bc from ChungLin_Tang@mentor.com ; Wed, 17 Aug 2016 05:17:38 -0700 Received: from [0.0.0.0] (147.34.91.1) by svr-orw-fem-03.mgc.mentorg.com (147.34.97.39) with Microsoft SMTP Server id 14.3.224.2; Wed, 17 Aug 2016 05:17:37 -0700 Subject: Re: [patch, OpenACC] Fix reduction lowering segfault in omp-low To: Jakub Jelinek References: <7d98ae79-3cb7-03f2-bbd5-34a78d30dfd4@codesourcery.com> <20160815095745.GM14857@tucnak.redhat.com> CC: gcc-patches , Cesar Philippidis From: Chung-Lin Tang Message-ID: <1658f83f-e37e-ca90-df15-d8046f597090@codesourcery.com> Date: Wed, 17 Aug 2016 20:17:36 +0800 User-Agent: Mozilla/5.0 (Macintosh; Intel Mac OS X 10.11; rv:45.0) Gecko/20100101 Thunderbird/45.2.0 MIME-Version: 1.0 In-Reply-To: <20160815095745.GM14857@tucnak.redhat.com> X-IsSubscribed: yes On 2016/8/15 5:57 PM, Jakub Jelinek wrote: > On Mon, Aug 15, 2016 at 05:52:29PM +0800, Chung-Lin Tang wrote: >> Hi Jakub, >> This patch fixes an OpenACC reduction lowering segfault which >> triggers when nested acc loop directives are present. >> Cesar has reviewed this patch internally (since he mostly wrote >> the code originally) >> >> Patch has been tested and committed to gomp-4_0-branch, >> is this also okay for trunk? >> >> Thanks, >> Chung-Lin >> >> 2016-08-15 Chung-Lin Tang >> >> * omp-low.c (lower_oacc_reductions): Adjust variable lookup to use >> maybe_lookup_decl, to handle nested acc loop directives. > > Is this covered by an existing testcase in the testsuite? > If not, can you please add a testcase for it. > Otherwise LGTM (not extra happy about accepting any kinds of contexts, > but I hope the nesting diagnostics error out on OpenMP contexts mixed with > OpenACC ones and hope that there can't be some other OpenACC context around > that you wouldn't want to handle). Thanks, I've committed the patch along with a new testcase (full patch as attached). Testcase is also backported to gomp-4_0-branch. Thanks, Chung-Lin Index: omp-low.c =================================================================== --- omp-low.c (revision 239529) +++ omp-low.c (working copy) @@ -5660,10 +5660,19 @@ lower_oacc_reductions (location_t loc, tree clause outgoing = var; incoming = omp_reduction_init_op (loc, rcode, type); } - else if (ctx->outer) - incoming = outgoing = lookup_decl (orig, ctx->outer); else - incoming = outgoing = orig; + { + /* Try to look at enclosing contexts for reduction var, + use original if no mapping found. */ + tree t = NULL_TREE; + omp_context *c = ctx->outer; + while (c && !t) + { + t = maybe_lookup_decl (orig, c); + c = c->outer; + } + incoming = outgoing = (t ? t : orig); + } has_outer_reduction:; } Index: testsuite/c-c++-common/goacc/reduction-6.c =================================================================== --- testsuite/c-c++-common/goacc/reduction-6.c (revision 0) +++ testsuite/c-c++-common/goacc/reduction-6.c (revision 0) @@ -0,0 +1,58 @@ +/* Check if different occurences of the reduction clause on + OpenACC loop nests will compile. */ + +int foo (int N) +{ + int a = 0, b = 0, c = 0, d = 0, e = 0; + + #pragma acc parallel + { + #pragma acc loop + for (int i = 0; i < N; i++) + { + #pragma acc loop reduction(+:a) + for (int j = 0; j < N; j++) + a += 1; + } + } + + #pragma acc parallel + { + #pragma acc loop reduction(+:b) + for (int i = 0; i < N; i++) + { + #pragma acc loop + for (int j = 0; j < N; j++) + b += 1; + } + } + + #pragma acc parallel + { + #pragma acc loop reduction(+:c) + for (int i = 0; i < N; i++) + { + #pragma acc loop reduction(+:c) + for (int j = 0; j < N; j++) + c += 1; + } + } + + #pragma acc parallel loop + for (int i = 0; i < N; i++) + { + #pragma acc loop reduction(+:d) + for (int j = 0; j < N; j++) + d += 1; + } + + #pragma acc parallel loop reduction(+:e) + for (int i = 0; i < N; i++) + { + #pragma acc loop reduction(+:e) + for (int j = 0; j < N; j++) + e += 1; + } + + return a + b + c + d + e; +}