[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:
> 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) {
...
}
```



================
Comment at: openmp/libomptarget/src/omptarget.cpp:410
+        return OFFLOAD_FAIL;
+    }
 
----------------
jdenny wrote:
> jdoerfert wrote:
> > Can we print the error in non debug mode as well? Is there a way we can provide more information on the location and variable?
> I would prefer that too.  When I investigated this originally, I didn't find a good precedent in this runtime implementation, so I decided to be conservative.  Is there an existing diagnostic I should model this after, or is this a unique case?
Unfortunately, libomptarget does not have access to source code location. `__tgt_` API functions do not have a related argument, like the `ident_t *loc` argument of `__kmpc_` API functions in libomp. Of course, if there is demand for such a feature we can implement it, although we would need to change *EVERY* API function in libomptarget.


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

https://reviews.llvm.org/D83062