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 >