[PATCH, 4/16] Implement -foffload-alias

[Tom de Vries]

On 27/11/15 12:42, Tom de Vries wrote:
> On 23/11/15 12:41, Richard Biener wrote:
>> On Sat, 21 Nov 2015, Tom de Vries wrote:
>>
>>> >On 13/11/15 12:39, Jakub Jelinek wrote:
>>>> > >On Fri, Nov 13, 2015 at 12:29:51PM +0100, Richard Biener wrote:
>>>>>> > > > >thanks for the explanation. Filed as PR68331 - '[meta-bug]
>>>>>> fipa-pta
>>>>>> > > > >issues'.
>>>>>> > > > >
>>>>>> > > > >Any feedback on the '#pragma GCC
>>>>>> offload-alias=<none|pointer|all>' bit
>>>>>> > > > >above?
>>>>>> > > > >Is that sort of what you had in mind?
>>>>> > > >
>>>>> > > >Yes.  Whether that makes sense is another question of course.
>>>>> You can
>>>>> > > >annotate memory references with MR_DEPENDENCE_BASE/CLIQUE
>>>>> yourself
>>>>> > > >as well if you know dependences without the users intervention.
>>>> > >
>>>> > >I really don't like even the GCC offload-alias, I just don't see
>>>> anything
>>>> > >special on the offload code.  Not to mention that the same issue
>>>> is already
>>>> > >with other outlined functions, like OpenMP tasks or parallel
>>>> regions, those
>>>> > >aren't offloaded, yet they can suffer from worse alias/points-to
>>>> analysis
>>>> > >too.
>>> >
>>> >AFAIU there is one aspect that is different for offloaded code: the
>>> setup of
>>> >the data on the device.
>>> >
>>> >Consider this example:
>>> >...
>>> >unsigned int a[N];
>>> >unsigned int b[N];
>>> >unsigned int c[N];
>>> >
>>> >int
>>> >main (void)
>>> >{
>>> >   ...
>>> >
>>> >#pragma acc kernels copyin (a) copyin (b) copyout (c)
>>> >   {
>>> >     for (COUNTERTYPE ii = 0; ii < N; ii++)
>>> >       c[ii] = a[ii] + b[ii];
>>> >   }
>>> >
>>> >   ...
>>> >...
>>> >
>>> >At gimple level, we have:
>>> >...
>>> >#pragma omp target oacc_kernels \
>>> >   map(force_from:c [len: 2097152]) \
>>> >   map(force_to:b [len: 2097152]) \
>>> >   map(force_to:a [len: 2097152])
>>> >...
>>> >
>>> >[ The meaning of the force_from/force_to mappings is given in
>>> >include/gomp-constants.h:
>>> >...
>>> >     /* Allocate.  */
>>> >     GOMP_MAP_FORCE_ALLOC = (GOMP_MAP_FLAG_FORCE | GOMP_MAP_ALLOC),
>>> >     /* ..., and copy to device.  */
>>> >     GOMP_MAP_FORCE_TO = (GOMP_MAP_FLAG_FORCE | GOMP_MAP_TO),
>>> >     /* ..., and copy from device.  */
>>> >     GOMP_MAP_FORCE_FROM = (GOMP_MAP_FLAG_FORCE | GOMP_MAP_FROM),
>>> >     /* ..., and copy to and from device.  */
>>> >     GOMP_MAP_FORCE_TOFROM = (GOMP_MAP_FLAG_FORCE | GOMP_MAP_TOFROM),
>>> >...  ]
>>> >
>>> >So before calling the offloaded function, a separate alloc is done
>>> for a, b
>>> >and c, and the base pointers of the newly allocated objects are
>>> passed to the
>>> >offloaded function.
>>> >
>>> >This means we can mark those base pointers as restrict in the offloaded
>>> >function.
>>> >
>>> >Attached proof-of-concept patch implements that.
>>> >
>>>> > >We simply have some compiler internal interface between the
>>>> caller and
>>>> > >callee of the outlined regions, each interface in between those has
>>>> > >its own structure type used to communicate the info;
>>>> > >we can attach attributes on the fields, or some flags to indicate
>>>> some
>>>> > >properties interesting from aliasing POV.
>>>> > >We don't really need to perform
>>>> > >full IPA-PTA, perhaps it would be enough to a) record somewhere
>>>> in cgraph
>>>> > >the relationship in between such callers and callees (for
>>>> offloading regions
>>>> > >we already have "omp target entrypoint" attribute on the callee
>>>> and a
>>>> > >singler caller), tell LTO if possible not to split those into
>>>> different
>>>> > >partitions if easily possible, and then just for these pairs perform
>>>> > >aliasing/points-to analysis in the caller and the result record
>>>> using
>>>> > >cliques/special attributes/whatever to the callee side, so that
>>>> the callee
>>>> > >(outlined OpenMP/OpenACC/Cilk+ region) can then improve its alias
>>>> analysis.
>>> >
>>> >As a start, is the approach of this patch OK?
>> Works for me but leaving to Jakub to review for correctness.
>
> Attached patch is a complete version:
> - added ChangeLog
> - added missing function header comments
> - moved analysis to separate function
>    omp_target_base_pointers_restrict_p
> - added example in comment before analysis
> - fixed error in omp_target_base_pointers_restrict_p where I was using
>    GOMP_MAP_ALLOC but should have been using GOMP_MAP_FORCE_ALLOC
> - added testcases
>

This follow-up patch handles the case that we copy from/to pointers 
rather than declared variables:
...
        void foo (unsigned int *a, unsigned int *b)
        {
	 #pragma acc kernels copyout (a[0:2]) copyout (b[0:2])
	 {
	   a[0] = 0;
	   b[0] = 1;
	 }
        }
...

After gimplification, we have:
...
      foo (unsigned int * a, unsigned int * b)
      {
        unsigned int * b.0;
        unsigned int * a.1;

        b.0 = b;
        a.1 = a;
        #pragma omp target oacc_kernels \
	 map(force_from:*a.1 (*a) [len: 8]) \
	 map(alloc:a [pointer assign, bias: 0]) \
	 map(force_from:*b.0 (*b) [len: 8]) \
	 map(alloc:b [pointer assign, bias: 0])
        {
	 unsigned int * a.2;
	 unsigned int * b.3;

	 a.2 = a;
	 *a.2 = 0;
	 b.3 = b;
	 *b.3 = 1;
       }
      }
...

We don't bail out of omp_target_base_pointers_restrict_p when 
encountering 'map(alloc:a [pointer assign, bias: 0])', given that we can 
find the matching 'map(force_from:*a.1 (*a) [len: 8])'.

Using this and the previous patch, I'm able to do auto-parallelization 
on all the oacc kernels c test-cases, with the obvious exception of the 
testcases where some of used variables are mapped using the 'present' 
tag (in other words, missing the force tag).

Bootstrapped and reg-tested on x86_64.

OK for stage3 trunk?

Thanks,
- Tom

-------------- next part --------------
A non-text attachment was scrubbed...
Name: 0002-Handle-non-declared-variables-in-kernels-alias-analysis.patch
Type: text/x-patch
Size: 10410 bytes
Desc: not available
URL: <http://gcc.gnu.org/pipermail/gcc-patches/attachments/20151127/0638f0e7/attachment.bin>