Re: [Bulk] [OpenACC 0/7] host_data construct

[Jakub Jelinek <jakub at redhat dot com>]

On Mon, Nov 02, 2015 at 06:33:39PM +0000, Julian Brown wrote:
> Firstly, on trunk at least, use_device_ptr variables are restricted to
> pointer or array types: that restriction doesn't exist in OpenACC, nor
> actually could I find it in the OpenMP 4.1 document (my guess is the
> standards are supposed to match in this regard). I think that a program
> such as this should work:

So, after talking about this on omp-lang, it seems there is agreement
that only arrays and pointer types (or reference to arrays or pointers)
should be allowed in use_device_ptr clause and that for pointers/reference
to pointers it should probably act the way I've coded it up, i.e. that
for them it translates the pointer to point to corresponding object to the
one to which it points on the host.  It is too late to change the standard
now, but will be changed soon, and hopefully clarified in examples.

> void target_fn (int *targ_data);
> 
> int
> main (int argc, char *argv[])
> {
>   char out;
>   int myvar;
> #pragma omp target enter data map(to: myvar)
> 
> #pragma omp target data use_device_ptr(myvar) map(from:out)
>   {
>     target_fn (&myvar);
>     out = 5;
>   }
> 
>   return 0;
> }

That would make the above non-conforming for OpenMP.

> Secondly, attempts to use use_device_ptr on (e.g.
> dynamically-allocated) arrays accessed through a pointer cause an ICE
> with the existing trunk OpenMP code:
> 
> #include <stdlib.h>
> 
> void target_fn (char *targ_data);
> 
> int
> main (int argc, char *argv[])
> {
>   char *myarr, out;
> 
>   myarr = malloc (1024);
> 
> #pragma omp target data map(to: myarr[0:1024])
>   {
> #pragma omp target data use_device_ptr(myarr) map(from:out)
>     {
>       target_fn (myarr);
>       out = 5;
>     }
>   }
> 
>   return 0;
> }

Can't reproduce this ICE (at least not on gomp-4_5-branch, but there
aren't significant changes from the trunk there).

> Furthermore, this looks strange to me (006t.omplower):
> 
>   .omp_data_arr.5.out = &out;
>   myarr.8 = myarr;
>   .omp_data_arr.5.myarr = myarr.8;
>   #pragma omp target data map(from:out [len: 1]) use_device_ptr(myarr)
>     {
>       D.2436 = .omp_data_arr.5.myarr;
>       myarr = D.2436;
> 
> That's clobbering the original myarr variable, right?

Just use -fdump-tree-omplower-uid to see that it is a different variable.
Basically, for OpenMP use_device_ptr creates a private copy of the
pointer for the body of the target data construct, and that pointer is
assigned the target device's address.  For arrays the implementation
creates an artificial pointer variable (holding the start of the array
initially) and replaces all references to the array in the target data
body with dereference of the pointer.

	Jakub