[patch,openacc] use firstprivate pointers for subarrays in c and c++

[Cesar Philippidis]

Pointers are special in OpenACC. Depending on the context, they can
either be treated as a "scalar" or as special firstprivate pointer. This
is in contrast to OpenMP target pointers, which are always treated as
firstprivate pointers if I'm not mistaken. The difference between a
firstprivate scalar and pointer is that the contents of a firstprivate
scalar are preserved on the accelerator, whereas a firstprivate pointer
gets remapped by the runtime to point to an address in the device's
address space. Here are the rules for pointers that I worked out with
the ACC technical committee.

  1) pointers used in subarrays shall be treated as firstprivate
     pointers

  2) all other pointers are scalars

There is an exception to 2) when a pointer appears inside a data region.
E.g.

  #pragma acc data copy (ptr[0:100])
  {
    #pragma acc parallel loop
    for (i = ...)
      ptr[i] = ...
  }

Here the compiler should detect that ptr is nested inside an acc data
region, and add an implicit present(ptr[0:100]) clause to it, and not
present(ptr). Note that the implicit data clause rule only applies to
lexically scoped offloaded regions inside acc data regions. E.g.

  foo (int *ptr)
  {
    ...
    #pragma acc parallel loop
    for (i = ...)
       ptr[i] = ...
  }

  bar ()
  {
    ...
    #pragma acc data copy (ptr[0:100])
    {
      foo (ptr);
    }
  }

will result in an implicit firstprivate(ptr) clause of the scalar
variety, not a firstprivate_pointer(ptr).

The attached patch updates gcc to implement this behavior. Currently,
gcc treats all pointers as scalars in OpenACC. So, if you have a
subarray involving a data mapping pcopy(p[5:10]), the gimplifier would
translate this clause into

  map(tofrom:*(p + 3) [len: 5]) map(alloc:p [pointer assign, bias: 3])

The alloc pointer map is a problem, especially with subarrays, because
it effectively breaks all of the acc_* library functions involving
subarrays. This patch changes that alloc map clause into a
map(firstprivate:c [pointer assign, bias: 3]).

This patch also corrects behavior of the acc_* library functions when
they deal with shared-memory targets. Before, most of those libraries
were incorrectly trying to create data maps for shared-memory targets.
The new behavior is to propagate the the host pointers where applicable,
bypassing the data map altogether.

Since I had to change so many existing compiler test cases, I also took
the liberty to make some of warning and error messages generated by the
c and c++ front ends more specific to OpenACC. In the c++ front end, I
went one step further and added preliminary checking for duplicate
reference-typed data mappings. This check is not that exhaustive though,
but I did include a test case for OpenMP.

It should be noted that I still need to update the behavior of subarray
pointers in fortran. I'm just waiting until for the OpenMP 4.5 fortran
changes to land in trunk first.

Is this patch OK for trunk?

Cesar

-------------- next part --------------
A non-text attachment was scrubbed...
Name: firstprivate_subarrays-20160510.diff
Type: text/x-patch
Size: 61568 bytes
Desc: not available
URL: <http://gcc.gnu.org/pipermail/gcc-patches/attachments/20160510/07cb2f54/attachment.bin>