From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: (qmail 85886 invoked by alias); 8 Apr 2016 07:41:03 -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 85326 invoked by uid 89); 8 Apr 2016 07:41:02 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-1.9 required=5.0 tests=BAYES_00,KAM_LAZY_DOMAIN_SECURITY,RP_MATCHES_RCVD,SPF_HELO_PASS autolearn=ham version=3.3.2 spammy=sum X-HELO: mx1.redhat.com Received: from mx1.redhat.com (HELO mx1.redhat.com) (209.132.183.28) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with (AES256-GCM-SHA384 encrypted) ESMTPS; Fri, 08 Apr 2016 07:41:01 +0000 Received: from int-mx13.intmail.prod.int.phx2.redhat.com (int-mx13.intmail.prod.int.phx2.redhat.com [10.5.11.26]) (using TLSv1.2 with cipher ECDHE-RSA-AES256-GCM-SHA384 (256/256 bits)) (No client certificate requested) by mx1.redhat.com (Postfix) with ESMTPS id 157A285366; Fri, 8 Apr 2016 07:41:00 +0000 (UTC) Received: from tucnak.zalov.cz (ovpn-113-22.phx2.redhat.com [10.3.113.22]) by int-mx13.intmail.prod.int.phx2.redhat.com (8.14.4/8.14.4) with ESMTP id u387ewd1025073 (version=TLSv1/SSLv3 cipher=DHE-RSA-AES256-GCM-SHA384 bits=256 verify=NO); Fri, 8 Apr 2016 03:40:59 -0400 Received: from tucnak.zalov.cz (localhost [127.0.0.1]) by tucnak.zalov.cz (8.15.2/8.15.2) with ESMTP id u387euXC017046; Fri, 8 Apr 2016 09:40:56 +0200 Received: (from jakub@localhost) by tucnak.zalov.cz (8.15.2/8.15.2/Submit) id u387esiv017045; Fri, 8 Apr 2016 09:40:54 +0200 Date: Fri, 08 Apr 2016 07:41:00 -0000 From: Jakub Jelinek To: Cesar Philippidis Cc: "gcc-patches@gcc.gnu.org" , Nathan Sidwell Subject: Re: openacc reference reductions Message-ID: <20160408074054.GN19207@tucnak.redhat.com> Reply-To: Jakub Jelinek References: <56BA0257.6050607@codesourcery.com> <56BA06C3.90606@acm.org> <56BA10FC.90705@codesourcery.com> <56CB2A76.3090809@codesourcery.com> <57046C2B.6080002@codesourcery.com> <20160406142340.GZ19207@tucnak.redhat.com> <57056FCA.7040602@codesourcery.com> <20160407095657.GD19207@tucnak.redhat.com> <570734E3.7030601@codesourcery.com> MIME-Version: 1.0 Content-Type: text/plain; charset=us-ascii Content-Disposition: inline In-Reply-To: <570734E3.7030601@codesourcery.com> User-Agent: Mutt/1.5.24 (2015-08-30) X-IsSubscribed: yes X-SW-Source: 2016-04/txt/msg00375.txt.bz2 On Thu, Apr 07, 2016 at 09:34:43PM -0700, Cesar Philippidis wrote: > --- a/gcc/gimplify.c > +++ b/gcc/gimplify.c > @@ -5802,7 +5802,8 @@ omp_add_variable (struct gimplify_omp_ctx *ctx, tree decl, unsigned int flags) > flags |= GOVD_SEEN; > > n = splay_tree_lookup (ctx->variables, (splay_tree_key)decl); > - if (n != NULL && (n->value & GOVD_DATA_SHARE_CLASS) != 0) > + if (n != NULL && (n->value & GOVD_DATA_SHARE_CLASS) != 0 > + && ctx->region_type != ORT_ACC_PARALLEL) > { > /* We shouldn't be re-adding the decl with the same data > sharing class. */ Why? > @@ -6557,6 +6558,24 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, > break; > } > > +/* OpenACC parallel reductions need a present_or_copy clause to ensure > + that the original variable used in the reduction gets updated on > + the host. Scan the list of clauses for reduction so that any existing > + data clause can be adjusted if necessary. */ > + if (region_type == ORT_ACC_PARALLEL) > + { > + for (c = *list_p; c; c = OMP_CLAUSE_CHAIN (c)) > + { > + tree decl = NULL_TREE; > + > + if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_REDUCTION) > + continue; > + > + decl = OMP_CLAUSE_DECL (c); > + omp_add_variable (ctx, decl, GOVD_REDUCTION); > + } > + } > + And this looks also wrong, why? If I try under the debugger 3 cases: void f1 (int sum) { #pragma acc parallel reduction(+:sum) present_or_copy(sum) ; } void f2 (int sum) { #pragma acc parallel present_or_copy(sum) ; } void f3 (int sum) { #pragma acc parallel reduction(+:sum) ; } then I see the loop that starts with the while below doing the right thing already. In the first case you end up with GOVD_SEEN | GOVD_EXPLICIT | GOVD_REDUCTION | GOVD_MAP in the second with GOVD_SEEN | GOVD_EXPLICIT | GOVD_MAP and third one with GOVD_SEEN | GOVD_EXPLICIT | GOVD_REDUCTION That is where you IMHO should stop at the gimplify_scan_omp_clauses side, so don't modify neither omp_add_variable nor gimplify_scan_omp_clauses at all, and do everything else in gimplify_adjust_omp_clauses. That function walks the explicit clauses and has all the info gathered during gimplify_scan_omp_clauses available in the splay tree. So, you can do all the checking there. Say on OMP_CLAUSE_REDUCTION for the ORT_ACC_PARALLEL check the flags if they include GOVD_PRIVATE or GOVD_FIRSTPRIVATE, if yes, complain. Also check if GOVD_MAP is included, if not, add the extra OMP_CLAUSE_MAP tofrom. And, on OMP_CLAUSE_MAP, check if GOVD_REDUCTION is set on ORT_ACC_PARALLEL, and if yes, check if it is tofrom and complain otherwise. > while ((c = *list_p) != NULL) > { > bool remove = false; > @@ -6808,6 +6827,31 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, > > case OMP_CLAUSE_MAP: > decl = OMP_CLAUSE_DECL (c); > + if (region_type == ORT_ACC_PARALLEL) > + { > + tree t = DECL_P (decl) ? decl : TREE_OPERAND (decl, 0); > + splay_tree_node n = NULL; > + > + if (DECL_P (t)) > + n = splay_tree_lookup (ctx->variables, (splay_tree_key)t); > + > + if (n) > + { > + int kind = OMP_CLAUSE_MAP_KIND (c); > + > + OMP_CLAUSE_MAP_IN_REDUCTION(c) = 1; > + if ((kind & GOMP_MAP_TOFROM) != GOMP_MAP_TOFROM > + && kind != GOMP_MAP_FORCE_PRESENT > + && kind != GOMP_MAP_POINTER) > + { > + warning_at (OMP_CLAUSE_LOCATION (c), 0, > + "incompatible data clause with reduction " > + "on %qE; promoting to present_or_copy", > + DECL_NAME (t)); > + OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_TOFROM); > + } > + } > + } > if (error_operand_p (decl)) > remove = true; > switch (code) So the above is also wrong IMNSHO. Jakub