On Wed, Oct 21, 2015 at 9:00 PM, Nathan Sidwell <nat...@acm.org> 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.
Richard.
> nathan
