On Thu, Oct 22, 2015 at 9:48 AM, Richard Biener
<richard.guent...@gmail.com> wrote:
> 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.
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
