From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: (qmail 5686 invoked by alias); 22 Oct 2015 07:49:35 -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 5669 invoked by uid 89); 22 Oct 2015 07:49:33 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-2.0 required=5.0 tests=AWL,BAYES_00,FREEMAIL_FROM,RCVD_IN_DNSWL_LOW,SPF_PASS autolearn=ham version=3.3.2 X-HELO: mail-yk0-f171.google.com Received: from mail-yk0-f171.google.com (HELO mail-yk0-f171.google.com) (209.85.160.171) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with (AES128-GCM-SHA256 encrypted) ESMTPS; Thu, 22 Oct 2015 07:49:32 +0000 Received: by ykaz22 with SMTP id z22so75021082yka.2 for ; Thu, 22 Oct 2015 00:49:30 -0700 (PDT) MIME-Version: 1.0 X-Received: by 10.129.49.211 with SMTP id x202mr10853685ywx.147.1445500170004; Thu, 22 Oct 2015 00:49:30 -0700 (PDT) Received: by 10.37.117.136 with HTTP; Thu, 22 Oct 2015 00:49:29 -0700 (PDT) In-Reply-To: References: <5627DD78.9040302@acm.org> <5627E0DF.9050507@acm.org> Date: Thu, 22 Oct 2015 07:55:00 -0000 Message-ID: Subject: Re: [OpenACC 1/11] UNIQUE internal function From: Richard Biener To: Nathan Sidwell Cc: GCC Patches , Jakub Jelinek , Bernd Schmidt , Jason Merrill , "Joseph S. Myers" Content-Type: text/plain; charset=UTF-8 X-IsSubscribed: yes X-SW-Source: 2015-10/txt/msg02199.txt.bz2 On Thu, Oct 22, 2015 at 9:48 AM, Richard Biener wrote: > On Wed, Oct 21, 2015 at 9:00 PM, Nathan Sidwell wrote: >> This patch implements a new internal function that has a 'uniqueness' >> property. Jump-threading cannot clone it and tail-merging cannot combine >> multiple instances. >> >> The uniqueness is implemented by a new gimple fn, >> gimple_call_internal_unique_p. Routines that check for identical or >> cloneable calls are augmented to check this property. These are: >> >> * tree-ssa-threadedge, which is figuring out if jump threading is a win. >> Jump threading is inhibited. >> >> * gimple_call_same_target_p, used for tail merging and similar transforms. >> Two calls of IFN_UNIQUE will never be the same target. >> >> * tracer.c, which is determining whether to clone a region. >> >> Interestingly jump threading avoids cloning volatile asms (which it admits >> is conservatively safe), but the tracer does not. I wonder if there's a >> latent problem in tracer? >> >> The reason I needed a function with this property is to preserve the >> looping structure of a function's CFG. As mentioned in the intro, we mark >> up loops (using this builtin), so the example I gave has the following >> inserts: >> >> #pragma acc parallel ... >> { >> // single mode here >> #pragma acc loop ... >> IFN_UNIQUE (FORKING ...) >> for (i = 0; i < N; i++) // loop 1 >> ... // partitioned mode here >> IFN_UNIQUE (JOINING ...) >> >> if (expr) // single mode here >> #pragma acc loop ... >> IFN_UNIQUE (FORKING ...) >> for (i = 0; i < N; i++) // loop 2 >> ... // partitioned mode here >> IFN_UNIQUE (JOINING ...) >> } >> >> The properly nested loop property of the CFG is preserved through the >> compilation. This is important as (a) it allows later passes to reconstruct >> this looping structure and (b) hardware constraints require a partioned >> region end for all partitioned threads at a single instruction. >> >> Until I added this unique property, original bring-up of partitioned >> execution would hit cases of split loops ending in multiple cloned JOINING >> markers and similar cases. >> >> To distinguish different uses of the UNIQUE function, I use the first >> argument, which is expected to be an INTEGER_CST. I figured this better >> than using multiple new internal fns, all with the unique property, as the >> latter would need (at least) a range check in gimple_call_internal_unique_p >> rather than a simple equality. >> >> Jakub, IYR I originally had IFN_FORK and IFN_JOIN as such distinct internal >> fns. This replaces that scheme. >> >> ok? > > Hmm, I'd just have used gimple_has_volatile_ops on the call? That > should have the > desired effects. That is, whatever new IFNs you need are ok, but special-casing them is not necessary if you properly mark the calls as volatile. Richard. > Richard. > >> nathan