oacc kernels directive -- reductions

[Tom de Vries]

Hi,

in this email I'm trying to explain in detail what problem I'm running into with 
reductions in oacc kernels region, and how I think it could be solved.

Any advice is welcome.


OVERALL PROBLEM

The overall problem I'm trying to solve is to implement the oacc kernels 
directive in gcc, reusing pass_parallelize_loops.


OACC KERNELS

The oacc kernels region is a region with a series of loop nests, which are 
intended to run on the accelerator. The compiler needs to offload each loop nest 
to the accelerator, in the way most optimal for the accelerator.


PASS_PARALLELIZE_LOOPS

The pass analyzes loops. If the loop iterations are independent, and it looks 
beneficial to parallelize the loop, the loop is transformed.

A copy of the loop is made, that deals with:
- small loop iterations for which the overhead of starting several threads will
   be too big, or
- fixup loop iterations that are left in case the number of iterations is not
   divisible by the parallelization factor.

The original loop is transformed:
- References of local variables are replaced with dereferences of a new
   variable, which are initialized at loop entry with the addresses of the
   original variables (eliminate_local_variables)
- copy loop-non-local variables to a structure, and replace references with
   loads from a pointer to another (similar) structure
   (seperate_decls_in_region)
- The loop is replaced with an GIMPLE_OMP_FOR (with and empty body) and
   GIMPLE_OMP_CONTINUE
- The loop region is enveloped with GIMPLE_OMP_PARALLEL and GIMPLE_OMP_RETURN
- the loop region is omp-expanded using omp_expand_local


STATUS

I've created an initial implementation in vries/oacc-kernels, on top of the 
gomp-4_0-branch.


GOMP-4_0-BRANCH

In the gomp-4_0-branch, the kernels directive is translated as a copy of the 
oacc parallels directive.  So, the following stages are done:
- pass_lower_omp/scan_omp:
   - scan directive body for variables.
   - build up omp_context datastructures.
   - declare struct with fields corresponding to scanned variables.
   - declare function with pointer to struct
- pass_lower_omp/lower_omp:
   - declare struct
   - assign values to struct fields
   - declare pointer to struct
   - rewrite body in terms of struct fields using pointer to struct.
- omp_expand:
   - build up omp_region data-structures
   - split off region in separate function
   - replace region with call to oacc runtime function while passing function
     pointer to split off function


VRIES/OACC-KERNELS

The current mechanism of offloading (compiling a function for a different 
architecture) is using the lto-streaming. The parloops pass is located after the 
lto-streaming point which is too late. OTOH, the parloops pass needs alias info, 
which is only available after pass_build_ealias. So a copy of the parloops pass 
specialized for oacc kernels has been added after pass_build_ealias (plus a 
couple of passes to compensate for moving the pass up in the pass list).

The new pass does not use the lowering (first 2 steps of loop transform) of 
parloops. The lowering is already done by pass_omp_lower.

The omp-expansion of the oacc-kernels region (done in gomp-4_0-branch) is 
skipped, to allow first the alias analysis to work on the scope of the intact 
function, and the new pass to do the omp-expansion.

So, the new pass:
- analyses the loop for dependences
- if independent, transforms the loop:
   - The loop is replaced with an GIMPLE_OMP_FOR (kind_oacc_loop, with an empty
     body) and GIMPLE_OMP_CONTINUE
   - The GIMPLE_OACC_KERNELS is replaced with GIMPLE_OACC_PARALLEL
   - the loop region is omp-expanded using omp_expand_local

The gotchas of the implementation are:
- no support for reductions, nested loops, more than one loop nest in
   kernels region
- the fixup/low-it-count loop copy is still generated _inside_ the split off
   function


PROBLEM WITH REDUCTIONS

In the vries/oacc-kernels implementation, the lowering of oacc kernels (in 
pass_lower_omp) is done before any loop analysis. For reductions, that's not 
possible anymore, since that would mean that detection of reductions comes after 
handling of reductions.

The problem we're running into here, is that:
- on one hand, the oacc lowering is done on high gimple (scopes still intact
   because GIMPLE_BINDs are still present, no bbs and cfgs, eh not expanded, no
   ssa),
- otoh, loop analysis is done on low ssa gimple (bbs, cfgs, ssa, no scopes, eh
   expanded)

The parloops pass is confronted with a similar problem.

AFAIU, ideal pass reuse for parloops would go something like this: on ssa, you 
do loop analysis. You then insert omp pragmas that indicate what transformations 
you want. Then you go back from ssa gimple to high gimple representation, and 
you run omp-lower and omp-expand to do the actual transformations.

Things have been solved like this in parloops: the lowering of omp-lower is not 
reused in parloops, but instead  a different (but similar) lowering has been 
added. What is reused, is the omp-expand. We don't go back to pre-ssa, but the 
omp-expand code has been adapted to handle ssa code. And the parloops pass 
removes the loop cfg part and substitutes it for a GIMPLE_OMP_FOR, as would be 
the case for omp for directives present in a source with omp directives.

An advantage that parloops has, is that its scope is limited to a single loop. 
OTOH, the implementation for oacc kernels has to cooperate with other oacc 
constructs, f.i. if an array is already present on the device due to an earlier 
construct.

Furthermore, the lowering styles are different.

In the omp-lower pass, oacc lowering is done like this: we load from a struct 
field a pointer (D.2158), from which we load the thread-local pointer 
instantiation (c.3D.2134):
...
   D.2158 = .omp_data_iD.2148->cD.2150;
   c.3D.2134 = *D.2158;
   D.2137 = c.3D.2134 + D.2136;
   *D.2137 = D.2144;
...

In the parloops pass, we simply load the thread-local pointer instantiation from 
a struct field:
...
   c.4_53 = .paral_data_load.9_56->c;

   ...

   _7 = c.4_53 + _5;
   *_7 = _14;
...


POSSIBLE SOLUTIONS

I can think of these possible solutions:

1. rewrite reduction analysis from parloops to work on high gimple.

2. rewrite omp-lowering to work on ssa gimple

3. try to reuse parloops lowering, and patch it up such that it coorporates with 
code generated for other oacc constructs.


At the moment, I'm looking at the last option.

By lowering the kernels directive and its clauses at omp-low, but the associated 
body at parloops, I get the following code (after parloops lowering) for a 
vector add loop:
...
   <bb 8>:
   .omp_data_arr.10.c = &c;
   .omp_data_arr.10.b = &b;
   .omp_data_arr.10.a = &a;
   #pragma acc kernels map(from:c [len: 8]) map(to:b [len: 8]) map(to:a [len: 
8]) [child fn: main._omp_fn.0 (.omp_data_arr.10, .omp_data_sizes.11, 
.omp_data_kinds.12)]

   <bb 9>:
   .omp_data_i_38 = &.omp_data_arr.10;
   c.5_40 = c;
   a.6_44 = a;
   b.7_49 = b;
...

I think that if I manage to replace the last 3 with this I could have the vector 
add loop working:
...
   c_p = .omp_data_i_38->c;
   a_p = .omp_data_i_38->a;
   b_p = .omp_data_i_38->b;
   c.5_40 = *c_p;
   a.6_44 = *a_p;
   b.7_49 = *b_p;
...

 From there on, I'll try to get reductions working in a similar fashion.

Thanks,
- Tom