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

[George Rokos via Phabricator via Openmp-commits]

grokos 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:
> 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.


================
Comment at: openmp/libomptarget/src/device.h:183
+                         bool HasCloseModifier = false,
+                         bool HasPresentModifier = false);
   void *getTgtPtrBegin(void *HstPtrBegin, int64_t Size);
----------------
jdenny wrote:
> I see no remaining `getOrAllocTgtPtr` calls that make use of these default arguments.  Any objection to removing them in this patch?
Sure, they can go.


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

https://reviews.llvm.org/D83062