From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: (qmail 105263 invoked by alias); 9 Nov 2015 13:59:20 -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 104568 invoked by uid 89); 9 Nov 2015 13:59:19 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-2.6 required=5.0 tests=BAYES_00,FREEMAIL_FROM,RCVD_IN_DNSWL_LOW,SPF_PASS autolearn=ham version=3.3.2 X-HELO: mail-qg0-f51.google.com Received: from mail-qg0-f51.google.com (HELO mail-qg0-f51.google.com) (209.85.192.51) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with (AES128-GCM-SHA256 encrypted) ESMTPS; Mon, 09 Nov 2015 13:59:18 +0000 Received: by qgeb1 with SMTP id b1so90227113qge.1 for ; Mon, 09 Nov 2015 05:59:16 -0800 (PST) X-Received: by 10.140.19.172 with SMTP id 41mr10353556qgh.89.1447077556442; Mon, 09 Nov 2015 05:59:16 -0800 (PST) Received: from ?IPv6:2601:181:c000:c497:a2a8:cdff:fe3e:b48? ([2601:181:c000:c497:a2a8:cdff:fe3e:b48]) by smtp.googlemail.com with ESMTPSA id c103sm4905407qgd.36.2015.11.09.05.59.15 (version=TLSv1.2 cipher=ECDHE-RSA-AES128-GCM-SHA256 bits=128/128); Mon, 09 Nov 2015 05:59:15 -0800 (PST) Subject: Re: OpenACC Firstprivate To: Jakub Jelinek References: <563E01A4.20607@acm.org> <20151109134619.GQ5675@tucnak.redhat.com> Cc: GCC Patches , Cesar Philippidis From: Nathan Sidwell Message-ID: <5640A6B3.3030409@acm.org> Date: Mon, 09 Nov 2015 13:59:00 -0000 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:38.0) Gecko/20100101 Thunderbird/38.3.0 MIME-Version: 1.0 In-Reply-To: <20151109134619.GQ5675@tucnak.redhat.com> Content-Type: text/plain; charset=windows-1252; format=flowed Content-Transfer-Encoding: 7bit X-SW-Source: 2015-11/txt/msg00960.txt.bz2 On 11/09/15 08:46, Jakub Jelinek wrote: > On Sat, Nov 07, 2015 at 08:50:28AM -0500, Nathan Sidwell wrote: >> Index: gcc/gimplify.c >> =================================================================== > > If you want to switch to hexadecimal, you should change all values > in the enum to hexadecimal for consistency. ok. >> >> /* 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. I presumed it necessary, as it was on the branch. will remove. > >> @@ -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. Will investigate (later) > >> + >> + /* 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. Ok. It seemed (and it may become clearer with default handling added), that OpenACC and OpenMP scanned scopes in opposite orders. I remember trying to get the ACC code to scan in the same order, but came up blank. Anyway, you're right, it should say what OpenACC is trying. > 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), I don't think that can be avoided. The host doesn't have control over when the CTAs (a gang) start -- they may even be serialized onto the same physical HW. So each gang has to initialize its own instance. Or did you mean something else? > 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) > { > ... > } > } Hm, I suspect that is either ill formed or the std does not contemplate. > 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? I would have to investigate ... nathan