[Openmp-commits] [PATCH] D83062: [OpenMP] Implement TR8 `present` map type modifier in runtime (2/2)

[Joel E. Denny via Phabricator via Openmp-commits]

jdenny marked 2 inline comments as done.
jdenny added inline comments.


================
Comment at: openmp/libomptarget/src/device.cpp:208-211
+    } else if (HasPresentModifier) {
+      DP("Mapping required but does not exist%s for HstPtrBegin=" DPxMOD
+         ", Size=%ld\n",
+         (IsImplicit ? " (implicit)" : ""), DPxPTR(HstPtrBegin), Size);
----------------
jdenny wrote:
> grokos wrote:
> > jdenny wrote:
> > > grokos wrote:
> > > > I think this else-if should be moved right after the else-if that checks for explicit extension of mapping (i.e. after line 194) and outside the `else if (Size)` branch. If we have the present modifier then the data must be mapped already no matter whether a size is specified.
> > > In the case of `Size == 0`, as far as I can tell, nothing is ever placed in `HostDataToTargetMap`, with or without the patch under review.  Instead, `getOrAllocTgtPtr` always handles this case as already mapped and returns `NULL` for it.  I believe your suggested change would not affect this behavior.  However, when `HasPresentModifier`, your suggested change would cause the `Mapping required but does not exist` debug message to be printed for this case, incorrectly implying that this case is never already mapped and that a runtime error should follow.  A runtime error won't follow because `Size == 0`.
> > > 
> > > In the case of `Size != 0`, your suggested change would affect unified shared memory.  As far as I can tell, unified shared memory is handled similarly to the case of `Size == 0`: it's never placed in `HostDataToTargetMap`, and `getOrAllocTgtPtr` always handles it as already mapped.  However, when `HasPresentModifier`, your suggested change would cause `getOrAllocTgtPtr` to always handle unified shared memory as unmapped, and a runtime error will follow, so unified shared memory is then unusable with the `present` modifier.
> > > 
> > > In both cases, your suggested change seems to produce the wrong behavior.  Do you agree?
> > > 
> > > I need to add tests to cover these cases.
> > > In the case of Size == 0, as far as I can tell, nothing is ever placed in HostDataToTargetMap, with or without the patch under review. Instead, getOrAllocTgtPtr always handles this case as already mapped and returns NULL for it.
> > This is not always true. We can have zero size and still get a valid device address. Consider the example below:
> > ```
> > int *p = malloc(N);
> > #pragma omp target enter data map(to: p[0:N])
> > ...
> > #pragma omp target
> > {
> >   p[0] = 1;
> > }
> > ```
> > Here `p[0:N]` is mapped via `enter data` and address &p[0] is inserted into `HostDataToTargetMap`. When we encounter the target region, the compiler captures pointer `p` implicitly as it appears inside the target region, however the compiler doesn't know its size (it's a pointer, not a named array), therefore it will be mapped as p[0:0], i.e. as a zero-size pointer. `getOrAllocTgtPtr` will find address &p[0] in `HostDataToTargetMap` (`lr.isContained`) and return the corresponding device address. If we have demanded that `p[0:0]` is present on the device, then libomptarget must check for its presence even if size is 0. So if you invoke the target region like this:
> > ```
> > #pragma omp target map(present, alloc : p[0:0])
> > ```
> > then libomptarget must either:
> >   # Return the corresponding device address of `&p[0]` if `p[0:N]` has been mapped before or
> >   # Error out if `&p[0]` cannot be found in `HostDataToTargetMap` (e.g. if we had omitted the `enter data` directive in the example above) but we requested it to be present
> > The latter case is not covered by this version of the patch.
> > 
> > Regarding unified shared memory you're right. So my proposed change should be:
> > ```
> > if (contained) {
> > ...
> > } else if (explicit extension) {
> > ...
> > } else if (HasPresentModifier && (!USM || (USM && Close) ) {
> >   error out
> > } else if (Size) {
> > ...
> > }
> > ```
> > 
> Thanks for explaining.  What is the correct behavior in the following case?
> 
> ```
> #pragma omp target data map(alloc : p[0:0])
> #pragma omp target data map(present, alloc : p[0:0])
> ```
> Thanks for explaining. What is the correct behavior in the following case?
> 
> #pragma omp target data map(alloc : p[0:0])
> #pragma omp target data map(present, alloc : p[0:0])

My latest update assumes the second directive should produce a runtime error because the first directive doesn't actually map anything.  Please let me know if that's not right.

> So my proposed change should be:
> 
> if (contained) {
> ...
> } else if (explicit extension) {
> ...
> } else if (HasPresentModifier && (!USM || (USM && Close) ) {
>   error out
> } else if (Size) {
> ...
> }

I organized this a little differently to avoid repeating the USM conditions, but I intend it to be equivalent.



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

https://reviews.llvm.org/D83062