Re: [PATCH 0/5] OpenMP/PTX: improve correctness in SIMD regions

[Alexander Monakov <amonakov at ispras dot ru>]

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?

Thanks.
Alexander