Re: [OpenACC 0/7] host_data construct

[Jakub Jelinek <jakub@redhat.com>]

On Thu, Nov 19, 2015 at 02:26:50PM +0000, Julian Brown wrote:
> OK, thanks -- as to what the standard says, it's so ill-specified in
> this area that nothing can be learned about the behaviour of offloaded
> regions within host_data constructs, and my question about that on the
> technical mailing list is still unanswered (actually Nathan suggested
> in private mail that the conservative thing to do would be to disallow
> offloaded regions entirely within host_data constructs, so maybe that's
> the way to go).
> 
> OpenMP 4.5 seems to *not* specify the skipping-over behaviour for
> use_device_ptr variables (p105, lines 20-23):
> 
> "The is_device_ptr clause is used to indicate that a list item is a
> device pointer already in the device data environment and that it
> should be used directly. Support for device pointers created outside
> of OpenMP, specifically outside of the omp_target_alloc routine and the
> use_device_ptr clause, is implementation defined."
> 
> That suggests that use_device_ptr is a valid way to create device
> pointers for use in enclosed target regions: the behaviour I assumed
> was wrong for OpenACC. So I think my guess at the "most-obvious"
> behaviour was probably misguided anyway.

use_device_ptr kind of privatizes the variable, the private variable being
the device pointer corresponding to the host pointer outside of the target
data with use_device_ptr clause.

And, if you want to use that device pointer in a target region, it should be
on the is_device_ptr clause on the target construct.  See e.g.
libgomp.c/target-18.c testcase.
  int a[4];
...
  #pragma omp target data map(to:a)
  #pragma omp target data use_device_ptr(a) map(from:err)
  #pragma omp target is_device_ptr(a) private(i) map(from:err)
  {
    err = 0;
    for (i = 0; i < 4; i++)
      if (a[i] != 23 + i)
        err = 1;
  }
The implementation has this way a choice how to implement device pointers
(what use_device_ptr gives you, or say omp_target_alloc returns)
- either (GCC's choice at least for the XeonPhi and hopefully PTX, HSA does
not care, as it shares address space) implement them as host pointer
encoding the bits the target device wants to use, or some kind of
descriptor.  In the former case, is_device_ptr is essentially a
firstprivate, you bitwise copy the device pointer from the host to target
device, where you can dereference it etc.  In the descriptor case you'd
do some transformation of the host side representation of the device pointer
to the device side.

> 
> It's maybe even more complicated. Consider the example:
> 
> char x[1024];
> 
> #pragma acc enter data copyin(x)
> 
> #pragma acc host_data use_device(x)
> {
>   target_primitive(x);
>   #pragma acc parallel present(x)    [1]
>   {
>     x[5] = 0;                        [2]
>   }
> }

If it is unclear, I think disallowing acc {parallel,kernels} inside of
acc host_data might be too big hammer, but perhaps just erroring out
or warning during gimplification that if you (explicitly or implicitly)
try to map a var that is in use_device clause in some outer context,
it is either wrong, unsupported or will not do what users think?

I will double check on omp-lang, but supposedly we could for OpenMP
warn in similar cases (use_device_ptr clause instead of use_device),
except when it is passed to is_device_ptr clause, because I think the
behavior is just unspecified otherwise.
> 
> Here, the "present" clause marked [1] will fail (because 'x' is a
> target pointer now). If it's omitted, the array access [2] will cause an
> implicit present_or_copy to be used for the 'x' pointer (which again
> will fail, because now 'x' points to target data). Maybe what we
> actually need is,
> 
> #pragma acc host_data use_device(x)
> {
>   target_primitive(x);
>   #pragma acc parallel deviceptr(x)
>   {
>     ...
>   }
> }
> 
> with the deviceptr(x) clause magically substituted in the parallel
> construct, but I'm struggling to see how we could justify doing that
> when that behaviour's not mentioned in the spec at all.

Is deviceptr as above meant to work?  That is the OpenACC counterpart
of is_device_ptr, right?  If yes, then I'd suggest just warning if you
try to implicitly or explicitly map something use_device in outer contexts,
and just make sure you don't ICE on the cases where you warn.
If the standard does not say what it means, then it is unspecified
behavior...

	Jakub