On Mon, 10 Sep 2018 22:22:15 +0100
Jason Merrill <ja...@redhat.com> wrote:
> On Mon, Sep 10, 2018 at 7:07 PM, Julian Brown
> <jul...@codesourcery.com> wrote:
> > I think it's more accurate to say that OpenACC says nothing about
> > C++ references at all, nor about how unadorned pointers are mapped
> > in copy/copyin/copyout clauses. So arguably we get to choose
> > whatever we want, preferably based on the principle of least
> > surprise. (ICE'ing definitely counts as a surprise!)
> >
> > As noted in a previous email, PGI seems to treat pointers to
> > aggregates specially, mapping them as ptr[0:1], but it's unclear if
> > the same is true for pointers to scalars with their compiler.
> > Neither behaviour seems to be standard-mandated, but this patch
> > extends the idea to references to scalars nonetheless.
>
> That certainly seems like the most sensible way of handling references
> to non-arrays. [...]
To try to clarify things for myself a bit, I tried to figure out better
what the current OpenMP behaviour in GCC is, and what the equivalent
OpenACC behaviour should be. I think the handling of references can and
should match between the two APIs (though implementation details of the
patch to make that so need a little work still).
Pointers (without array sections) are a little more awkward: going by
what OpenMP 4.5 and OpenACC 2.5 say, there does seem to be a deliberate
difference in mapping behaviour, at least for cases that are specified.
Previously, I was confusing the cases marked (*) and (**) below a
little. So, we have:
== OpenMP 4.5 =====================================================
#include <stdio.h>
int
main (int argc, char* argv[])
{
int arr[32];
int &myref = arr[16];
int *myptr = &arr[18];
const char *sep = "";
for (int i = 0; i < 32; i++)
arr[i] = i;
//#pragma omp target // mapped as firstprivate: no effect on host
//#pragma omp target defaultmap(tofrom:scalar) // works
#pragma omp target map(tofrom:myref) // works
{
myref = 1000;
}
#pragma omp target enter data map(to:arr[0:32])
//#pragma omp target // works, mapped as zero-length array section (*)
//#pragma omp target map(tofrom:myptr) // crashes (**)
#pragma omp target map(tofrom:myptr[0:1]) // works
{
*myptr = 2000;
}
#pragma omp target exit data map(from:arr[0:32])
for (int i = 0; i < 32; i++, sep = ", ")
printf ("%s%d", sep, arr[i]);
printf ("\n");
return 0;
}
== OpenACC 2.5 ====================================================
#include <stdio.h>
int
main (int argc, char* argv[])
{
int arr[32];
int &myref = arr[16];
int *myptr = &arr[18];
const char *sep = "";
for (int i = 0; i < 32; i++)
arr[i] = i;
//#pragma acc parallel // mapped as firstprivate: no effect on host
#pragma acc parallel copy(myref) // works
{
myref = 1000;
}
#pragma acc enter data copyin(arr[0:32])
//#pragma acc parallel // crashes (*)
//#pragma acc parallel copy(myptr) // crashes (**)
//#pragma acc parallel copy(myptr[0:1]) // works
//#pragma acc parallel present(myptr) // runtime error, not present
#pragma acc parallel present(myptr[0:1]) // works
{
*myptr = 2000;
}
#pragma acc exit data copyout(arr[0:32])
for (int i = 0; i < 32; i++, sep = ", ")
printf ("%s%d", sep, arr[i]);
printf ("\n");
return 0;
}
===================================================================
The pointer-mapping cases marked (*), implicit mapping, are the ones
specified in OpenMP 4.5 to map as zero-length array sections. For
OpenACC the pointer is considered a scalar so is mapped as bits (so the
host pointer causes the target to crash on dereference).
The cases marked (**) -- also maybe applicable to C++ "this" --
currently copy as bits on OpenMP and on OpenACC, but could be changed
to map like length-one array sections. Or, they could raise a warning.
There's no apparent difference between OpenMP and OpenACC there though
(in specified behaviour and/or implementation? Despite what I thought
previously) so that's probably a decision for another day.
Cheers,
Julian
