Re: [OpenACC 1/11] UNIQUE internal function

[Richard Biener <richard dot guenther at gmail dot com>]

On Thu, Oct 22, 2015 at 9:48 AM, Richard Biener
<richard.guenther@gmail.com> wrote:
> On Wed, Oct 21, 2015 at 9:00 PM, Nathan Sidwell <nathan@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