[PATCH] D65835: [OpenMP] Fix map/is_device_ptr with DSA on combined directive

[Hal Finkel via Phabricator via cfe-commits]

hfinkel added a comment.

In D65835#1623903 <https://reviews.llvm.org/D65835#1623903>, @ABataev wrote:

> In D65835#1623883 <https://reviews.llvm.org/D65835#1623883>, @hfinkel wrote:
>
> > In D65835#1623814 <https://reviews.llvm.org/D65835#1623814>, @hfinkel wrote:
> >
> > > In D65835#1623772 <https://reviews.llvm.org/D65835#1623772>, @ABataev wrote:
> > >
> > > > In D65835#1623756 <https://reviews.llvm.org/D65835#1623756>, @kkwli0 wrote:
> > > >
> > > > > In D65835#1622042 <https://reviews.llvm.org/D65835#1622042>, @ABataev wrote:
> > > > >
> > > > > > > I want to be sure we're on the same page: For OpenMP 5.0, should we allow is_device_ptr with the private clauses?
> > > > > >
> > > > > > Yes, since it is allowed by the standard.
> > > > >
> > > > >
> > > > > Umm ... I probably missed some earlier discussions!  What would be the behavior of the following code?
> > > > >
> > > > >   p = omp_target_alloc(...);
> > > > >   #pragma omp target private(p) is_device_ptr(p)
> > > > >     p[...] = ...;   // crash or not?
> > > > >
> > > >
> > > >
> > > > It must crush, I assume. The main problem is that this construct is allowed by the standard.
> > >
> > >
> > > Yep. We should add a warning message for it.
> >
> >
> > Upon further reflection, this is not clearly allowed by the standard. My experience is that, when reading standards, sometimes things are disallowed by contradiction (i.e., the standard does not define some behavior, and what the standard does say that's relevant is self contradictory). In this case, 2.19.3 says that list items which are privatized (and which are used) undergo replacement (with new items created as specified) while 2.12.5 says that "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." A given list item cannot simultaneously be "used directly" (2.12.5) and also undergo replacement: "Inside the construct, all references to the original list item are replaced by references to a new list item received by the task or SIMD lane" (2.19.3). Thus, it may be disallowed.
>
>
> I think, this combination is still allowed. I assume that
>
>   #Pragma omp target parallel is_device_ptr(a) <dsa_clause>(a)
>
>
> is the same as
>
>   #pragma omp target is_device_ptr(a)
>   #pragma omp parallel <dsa_clause>(a)
>
>
> i.e. datasharing clauses are applied to inner sub-regions, not the target region itself.


With the parallel, that makes sense to me. In that case, however, I'd imagine that the privitization works as normal and the code wouldn't crash. Right?


Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D65835/new/

https://reviews.llvm.org/D65835