[1/3] OpenACC reductions

[Nathan Sidwell]

On 11/03/15 10:46, Jakub Jelinek wrote:
> On Mon, Nov 02, 2015 at 11:18:37AM -0500, Nathan Sidwell wrote:
>> This is the core execution bits of OpenACC reductions.
>>
>> We have a new internal fn 'IFN_GOACC_REDUCTION' and a new target hook
>> goacc.reduction, to lower it on the target compiler.
>
> So, let me start with a few questions:
> 1) does OpenACC allow UDRs or only the built-in reductions?  If it
>     does not allow UDRs, do you have it covered by testcases that you
>     disallow parsing of them (e.g. when you have

no UDR reductions.  Will check test cases for that.

> #pragma omp declare reduction (xyz: struct S: omp_out.x += omp_in.y) initializer (omp_priv = { 5 })
> #pragma acc parallel reduction (xyz: var_with_type_S)
>     )?

> 2) how do you expand the reductions in the end when targetting host fallback
>     or when targetting non-PTX targets?

That's what default_goacc_reduction is doing.

(I see its comment hasn't caught up with the changes I made during the merge. 
Will fix)

    LHS-opt = IFN_RED (KIND, RES_PTR, VAR, LEVEL, OP, OFFSET)
    If RES_PTR is not integer-zerop:
        SETUP - emit 'LHS = *RES_PTR', LHS = NULL
        TEARDOWN - emit '*RES_PTR = VAR'
    If LHS is not NULL
        emit 'LHS = VAR'

This is the correct behaviour for a single-threaded  loop.  Of course the loop 
could go on to be parallelized in the normal way -- or additional conversion to 
openmp constructs along the same lines as we discussed for the GOACC_LOOP function.

Does that help?

nathan