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

Tue Jul 7 17:07:31 PDT 2020

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);
----------------
grokos wrote:
> jdenny wrote:
> > grokos wrote:
> > > jdenny wrote:
> > > > 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.
> > > > 
> > > > 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])
> > > 
> > > For the first one, if the object `p` is pointing to has been mapped before then there is no noticeable effect, the only thing that happens is that `getOrAllocTgtPtr` increments the RefCount of the mapping and returns the corresponding device address. If the object hasn't been mapped before and because size is zero, `getOrAllocTgtPtr` will return NULL (there is nothing to be allocated).
> > > 
> > > With the `present` modifier, if the object is already mapped, `getOrAllocTgtPtr` increments the RefCount and returns the corresponding device address (just like in the former case). If the object is not mapped, the `present` modifier will trigger an error - in other words `NULL` is not an option if we have requested something to be present.
> > Sorry, I meant that these directives are nested and the array was not previously mapped.  It sounds like the outer directive does not map it in that case, and so the inner directive produces a runtime error.
> Correct.
Thanks for all the careful explanations!  I believe the logic is now correct for zero-length array sections.

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

https://reviews.llvm.org/D83062