oacc kernels directive -- reductions

[Richard Biener]

On Tue, 14 Oct 2014, Tom de Vries wrote:

> 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.

Not really understanding how 3) can replace 1) or 2) ... but what
I understand is that OACC lowering happens at pass_lower_omp
(no CFG or loops or SSA).  If we want to keep that the "proper"
choice of dealing with this high-level OACC "kernel" directive
is to pass it down somehow and deal with it later when loops + SSA
are available.  After all you _are_ re-using parloops for the
"lowering".

So - can't OACC kernel lowering at lower_omp time simply annotate loops?
Like in a way #pragma ivdep is handled?  Maybe this should even happen
from inside the parser (who knows where loops are)?  Then at CFG / loop
build time this information is transfered to loop meta-data (same
as with IVDEP) and a parloop pass somewhere in early opts can do
the right thing(TM) on the marked loops?

Thanks,
Richard.

> Thanks,
> - Tom
> 
> 

-- 
Richard Biener <rguenther@suse.de>
SUSE / SUSE Labs
SUSE LINUX Products GmbH - Nuernberg - AG Nuernberg - HRB 16746
GF: Jeff Hawn, Jennifer Guild, Felix Imend"orffer