[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 an inline comment 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:
> > > 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])
```


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

https://reviews.llvm.org/D83062