This is the mail archive of the gcc@gcc.gnu.org mailing list for the GCC project.


Index Nav: [Date Index] [Subject Index] [Author Index] [Thread Index]
Message Nav: [Date Prev] [Date Next] [Thread Prev] [Thread Next]
Other format: [Raw text]

oacc kernels directive -- reductions


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


Index Nav: [Date Index] [Subject Index] [Author Index] [Thread Index]
Message Nav: [Date Prev] [Date Next] [Thread Prev] [Thread Next]