This is the mail archive of the gcc@gcc.gnu.org mailing list for the GCC project.


Index Nav: [Date Index] [Subject Index] [Author Index] [Thread Index]
Message Nav: [Date Prev] [Date Next] [Thread Prev] [Thread Next]
Other format: [Raw text]

Re: [gomp4] Questions about "declare target" and "target update" pragmas


On Thu, Mar 19, 2015 at 05:49:47PM +0300, Ilya Verbin wrote:
> > > void foo (int a1[])
> > > {
> > >   #pragma omp target
> > >     {
> > >       a1[10]++;
> > >       a2[10]++;
> > >     }
> > > }
> > 
> > That is a buggy test.  int a1[] function argument is changed
> > into int *a1, so it is actually
> > #pragma omp target map(tofrom:a1, a2)
> 
> Actually, it copies only a1 pointer, since a2 points to the global array.

Sure, the implicit map(tofrom:a2) doesn't really do anything, since a2
is clearly already mapped and will stay to be mapped, and is a global var.
> 
> > {
> >   a1[10]++;
> >   a2[10]++;
> > }
> > which copies the a1 pointer to the device by value (no pointer
> > transformation).
> > Perhaps the testcase writer meant to use #pragma omp target map(a1[10])
> > instead (or map(a1[0:50])?
> 
> If I understand correctly, it's not allowed to map global target arrays this
> way, since it's already present in the initial device data environment:

It of course is allowed.  It just means that it doesn't allocate new memory
(sizeof(int) large in the map(a1[10]) case, sizeof(int)*50 large in the map(a1[0:50])
case), nor copy the bytes around, all it will do is allocate memory for the
target copy of the a1 pointer, and do pointer transformation such that the
a1 pointer on the target will point to the global target a1 array.
Without the map(a1[10]) or map(a1[0:50]) clauses (i.e. implicit map(tofrom:a1))
it does similar thing, except it copies the pointer value to the target (and
back at the end of the region) instead, which is not what you want...

> So, to fix this testcase I can just remove the "int a1[]" function argument, and
> add some "#pragma omp target update" where needed.

Well, supposedly the test was meant to test it with a pointer parameter,
otherwise why would there be both a1 and a2 when a2 would be enough?

	Jakub


Index Nav: [Date Index] [Subject Index] [Author Index] [Thread Index]
Message Nav: [Date Prev] [Date Next] [Thread Prev] [Thread Next]