From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: (qmail 2026 invoked by alias); 9 Nov 2015 13:46:29 -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 2014 invoked by uid 89); 9 Nov 2015 13:46:28 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-0.6 required=5.0 tests=AWL,BAYES_50,RP_MATCHES_RCVD,SPF_HELO_PASS autolearn=ham version=3.3.2 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; Mon, 09 Nov 2015 13:46:25 +0000 Received: from int-mx14.intmail.prod.int.phx2.redhat.com (int-mx14.intmail.prod.int.phx2.redhat.com [10.5.11.27]) by mx1.redhat.com (Postfix) with ESMTPS id 20B572D4; Mon, 9 Nov 2015 13:46:24 +0000 (UTC) Received: from tucnak.zalov.cz (ovpn-116-34.ams2.redhat.com [10.36.116.34]) by int-mx14.intmail.prod.int.phx2.redhat.com (8.14.4/8.14.4) with ESMTP id tA9DkMEN010449 (version=TLSv1/SSLv3 cipher=DHE-RSA-AES256-GCM-SHA384 bits=256 verify=NO); Mon, 9 Nov 2015 08:46:23 -0500 Received: from tucnak.zalov.cz (localhost [127.0.0.1]) by tucnak.zalov.cz (8.15.2/8.15.2) with ESMTP id tA9DkKM0025202; Mon, 9 Nov 2015 14:46:20 +0100 Received: (from jakub@localhost) by tucnak.zalov.cz (8.15.2/8.15.2/Submit) id tA9DkJH9025201; Mon, 9 Nov 2015 14:46:19 +0100 Date: Mon, 09 Nov 2015 13:46:00 -0000 From: Jakub Jelinek To: Nathan Sidwell Cc: GCC Patches , Cesar Philippidis Subject: Re: OpenACC Firstprivate Message-ID: <20151109134619.GQ5675@tucnak.redhat.com> Reply-To: Jakub Jelinek References: <563E01A4.20607@acm.org> MIME-Version: 1.0 Content-Type: text/plain; charset=us-ascii Content-Disposition: inline In-Reply-To: <563E01A4.20607@acm.org> User-Agent: Mutt/1.5.23 (2014-03-12) X-IsSubscribed: yes X-SW-Source: 2015-11/txt/msg00957.txt.bz2 On Sat, Nov 07, 2015 at 08:50:28AM -0500, Nathan Sidwell wrote: > Index: gcc/gimplify.c > =================================================================== > --- gcc/gimplify.c (revision 229892) > +++ gcc/gimplify.c (working copy) > @@ -108,9 +108,15 @@ enum omp_region_type > /* Data region with offloading. */ > ORT_TARGET = 32, > ORT_COMBINED_TARGET = 33, > + > + ORT_ACC = 0x40, /* An OpenACC region. */ > + ORT_ACC_DATA = ORT_ACC | ORT_TARGET_DATA, /* Data construct. */ > + ORT_ACC_PARALLEL = ORT_ACC | ORT_TARGET, /* Parallel construct */ > + ORT_ACC_KERNELS = ORT_ACC | ORT_TARGET | 0x80, /* Kernels construct. */ > + > /* Dummy OpenMP region, used to disable expansion of > DECL_VALUE_EXPRs in taskloop pre body. */ > - ORT_NONE = 64 > + ORT_NONE = 0x100 > }; If you want to switch to hexadecimal, you should change all values in the enum to hexadecimal for consistency. > > /* Gimplify hashtable helper. */ > @@ -377,6 +383,12 @@ new_omp_context (enum omp_region_type re > else > c->default_kind = OMP_CLAUSE_DEFAULT_UNSPECIFIED; > > + c->combined_loop = false; > + c->distribute = false; > + c->target_map_scalars_firstprivate = false; > + c->target_map_pointers_as_0len_arrays = false; > + c->target_firstprivatize_array_bases = false; Why this? c is XCNEW allocated, so zero initialized. > @@ -5667,11 +5682,13 @@ omp_add_variable (struct gimplify_omp_ct > /* We shouldn't be re-adding the decl with the same data > sharing class. */ > gcc_assert ((n->value & GOVD_DATA_SHARE_CLASS & flags) == 0); > - /* The only combination of data sharing classes we should see is > - FIRSTPRIVATE and LASTPRIVATE. */ > nflags = n->value | flags; > - gcc_assert ((nflags & GOVD_DATA_SHARE_CLASS) > - == (GOVD_FIRSTPRIVATE | GOVD_LASTPRIVATE) > + /* The only combination of data sharing classes we should see is > + FIRSTPRIVATE and LASTPRIVATE. However, OpenACC permits > + reduction variables to be used in data sharing clauses. */ > + gcc_assert ((ctx->region_type & ORT_ACC) != 0 > + || ((nflags & GOVD_DATA_SHARE_CLASS) > + == (GOVD_FIRSTPRIVATE | GOVD_LASTPRIVATE)) > || (flags & GOVD_DATA_SHARE_CLASS) == 0); Are you sure you want to give up on any kind of consistency checks for OpenACC? If only reduction is special on OpenACC, perhaps you could tweak the assert for that instead? Something that can be done incrementally of course. > + > + /* OpenMP doesn't look in outer contexts to find an > + enclosing data clause. */ I'm puzzled by the comment. OpenMP does look in outer context for clauses that need that (pretty much all closes but private), that is the do_outer: recursion in omp_notice_variable. Say for firstprivate in order to copy (or copy construct) the private variable one needs the access to the outer context's var etc.). So perhaps it would help to document what you are doing here for OpenACC and why. > + struct gimplify_omp_ctx *octx = ctx->outer_context; > + if ((ctx->region_type & ORT_ACC) && octx) > + { > + omp_notice_variable (octx, decl, in_code); > + > + for (; octx; octx = octx->outer_context) > + { > + if (!(octx->region_type & (ORT_TARGET_DATA | ORT_TARGET))) > + break; > + splay_tree_node n2 > + = splay_tree_lookup (octx->variables, > + (splay_tree_key) decl); > + if (n2) > + { > + nflags |= GOVD_MAP; > + goto found_outer; > + } > + } > } > - else if (nflags == flags) > - nflags |= GOVD_MAP; > + The main issue I have is with the omp-low.c changes. I see: "2.5.9 private clause The private clause is allowed on the parallel construct; it declares that a copy of each item on the list will be created for each parallel gang. 2.5.10 firstprivate clause The firstprivate clause is allowed on the parallel construct; it declares that a copy of each item on the list will be created for each parallel gang, and that the copy will be initialized with the value of that item on the host when the parallel construct is encountered." but looking at what you actually emit looks like standard present_copyin clause I think with a private variable defined in the region where the value of the present_copyin mapped variable is assigned to the private one. This I'm afraid performs often two copies rather than just one (one to copy the host value to the present_copyin mapped value, another one in the region), but more importantly, if the var is already mapped, you could initialize the private var with old data. Say int arr[64]; // initialize arr #pragma acc data copyin (arr) { // modify arr on the host # pragma acc parallel firstprivate (arr) { ... } } Is that really what you want? If not, any reason not to implement GOMP_MAP_FIRSTPRIVATE and GOMP_MAP_FIRSTPRIVATE_INT on the libgomp oacc-* side and just use the OpenMP firstprivate handling in omp-low.c? Jakub