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

[Kelvin Li via Phabricator via cfe-commits]

kkwli0 added a comment.

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.


That is what I thought.  Specifying these two clauses on the target construct creates ambiguity on which p it referred to inside the construct.  The private p or the pointer p that stores the device address?  I will work with the committee to fix the spec.


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D65835