[OpenACC 1/11] UNIQUE internal function

[Nathan Sidwell]

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?

nathan
-------------- next part --------------
A non-text attachment was scrubbed...
Name: 01-trunk-unique.patch
Type: text/x-patch
Size: 4820 bytes
Desc: not available
URL: <http://gcc.gnu.org/pipermail/gcc-patches/attachments/20151021/9682f6bd/attachment.bin>