On Wed, Jan 18, 2017 at 3:11 PM, Alexander Monakov <amona...@ispras.ru> wrote:
> On Wed, 18 Jan 2017, Richard Biener wrote:
>> > After OpenMP lowering, inlining might break this by inlining functions with
>> > address-taken locals into SIMD regions. For now, such inlining is
>> > disallowed
>> > (this penalizes only SIMT code), but eventually that can be handled by
>> > collecting those locals into an allocated struct in a similar manner.
>>
>> Can you do the allocation fully after inlining instead?
>
> Hm. I'm not really sure what you mean, because I may not emit GIMPLE that
> takes
> addresses of an incomplete struct's fields, and changing layout of an existing
> completely layed out struct is not trivial either. But I have an idea, see
> below.
>
> Let's consider what the last patch implements; starting from
>
> #pragma omp simd private(tmp)
> for (int i = n1; i < n2; i++)
> foo (&tmp);
>
> it emits GIMPLE that looks like this:
>
> struct {
> int tmp;
> } *omp_simt = IFN_GOMP_SIMT_ENTER (sizeof *omp_simt);
>
> /* This temporary is needed because we populate the struct and (re)gimplify
> references to privatized variables in one pass; replacing 'tmp' directly
> with '&omp_simt->tmp' wouldn't work, because struct layout is not known
> until all fields are added, and gimplification wouldn't be able to emit
> the corresponding MEM_REF. */
> int *tmp_ptr = &omp_simt->tmp;
>
> for (int i = n1; i < n2; i++)
> foo (tmp_ptr);
>
> *.omp_simt = {CLOBBER};
> IFN_GOMP_SIMT_EXIT (.omp_simt);
>
>
> So I guess a way to keep allocation layout implicit until after inlining is
> this: instead of exposing the helper struct in the IR immediately, somehow
> keep
> it on the side, associated only with the SIMT region, and not finalized. This
> would allow to populate it as needed during inlining, but the downside is that
> references to privatized vars would get weirder: they would need to be via
> IFNs
> that track association with the loop and the privatized variable. Like this:
>
> void *omp_simt = IFN_GOMP_SIMT_ENTER_BY_UID (simduid);
>
> int *tmp_ptr = IFN_GOMP_SIMT_VAR_REF (omp_simt, simduid, uid_for_tmp);
>
> for (...)
> foo (tmp_ptr);
>
> *tmp_ptr = {CLOBBER}; /* ??? for each privatized variable? */
> IFN_GOMP_SIMT_EXIT (.omp_simt);
>
> (note how in this scheme we'd need to emit separate CLOBBERs for each field)
>
> But absence of explicit struct would hurt alias analysis I'm afraid: it
> wouldn't
> be able to deduce that references to different privatized variable do not
> alias
> until after calls to SIMT_VAR_REF are replaced. Or is that not an issue?
It probably is.
But I guess I was asking whether you could initially emit
void *omp_simt = IFN_GOMP_SIMT_ENTER (0);
for (int i = n1; i < n2; i++)
foo (&tmp);
IFN_GOMP_SIMT_EXIT (omp_simt);
and only after inlining do liveness / use analysis of everything between
SIMT_ENTER and SIMT_EXIT doing the rewriting only at that point.
Richard.
> Thanks.
> Alexander
