[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 3 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:
> 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.


================
Comment at: openmp/libomptarget/src/omptarget.cpp:266-269
+      // TODO: Do we actually need to check the present modifier here?  I
+      // haven't found an example where the getOrAllocTgtPtr call here fails
+      // because, so far in my attempts, it's always already present due to an
+      // earlier map entry for the object holding the pointer.
----------------
grokos wrote:
> Correct. The only case we can have a `PTR_AND_OBJ` entry is if we have a pointer which is member of a struct or element of an array. This means that whenever we have something like:
> ```
> #pragma omp target map(s.p[0:N])
> ```
> clang will first map the struct/array and then we'll visit the `PTR_AND_OBJ` entry, so the space for the pointer itself has been allocated already. I think it's safe to skip this check here (or you can leave it as a form of defensive programming).
Thanks for confirming.  I'll leave it but adjust the comment not to be a todo.


================
Comment at: openmp/libomptarget/src/omptarget.cpp:401
         UpdateRef, IsHostPtr);
-    DP("There are %" PRId64 " bytes allocated at target address " DPxMOD
-        " - is%s last\n", data_size, DPxPTR(TgtPtrBegin),
-        (IsLast ? "" : " not"));
+    if (TgtPtrBegin) {
+      DP("There are %" PRId64 " bytes allocated at target address " DPxMOD
----------------
This should be `TgtPtrBegin || !data_size`, mirroring similar code in `target_data_begin`.


================
Comment at: openmp/libomptarget/src/omptarget.cpp:410
+        return OFFLOAD_FAIL;
+    }
 
----------------
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?


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D83062