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

[Cesar Philippidis <cesar_philippidis at mentor dot com>]

Ping.

Cesar

On 05/10/2016 01:29 PM, Cesar Philippidis wrote:
> 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
>