[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 5 inline comments as done.
jdenny added inline comments.


================
Comment at: openmp/libomptarget/src/omptarget.cpp:268
+      // 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.  For example:
+      //
----------------
grokos wrote:
> jdenny wrote:
> > grokos wrote:
> > > jdenny wrote:
> > > > This comment is incorrect.  We can also have PTR_AND_OBJ in the case of `omp declare target link` as discussed at <http://lists.llvm.org/pipermail/openmp-dev/2020-July/003586.html>.
> > > Right, the comment is not relevant anymore. But I think the conveyed message is still correct. Apart from structs/arrays, if we have a `PTR_AND_OBJ` for a global pointer then by definition the pointer itself has been mapped (since it's a global). Maybe update the comment? The logic is still correct I think.
> > Yes, I believe the code is right.  I'll try to generalize the comment.
> I was just reminded that globals marked with `declare target link` are also PTR_AND_OBJ entries. `declare target link` already works in the exact same way as the proposed change we discussed at the July 15 telecon.
If we're so sure it's impossible to have `!Pointer_TgtPtrBegin` here, should we have an assert?

That way, if the situation evolves, testing will hopefully reveal when our assumptions here are wrong.  Based on my understanding now, it would probably indicate a bug somewhere.  But I guess it might just mean the assert then needs to be removed and the comment needs to be updated.  Either way, it will be good to know about the change.

My understanding is that assert would only be active for debug builds.  For release builds, we could still have the `if` and return `OFFLOAD_FAIL` for the sake of defensive programming.


================
Comment at: openmp/libomptarget/src/omptarget.cpp:416
+                DPxPTR(HstPtrBegin), data_size);
+        return OFFLOAD_FAIL;
+      }
----------------
jdenny wrote:
> This is right for `omp target exit data`, as in `openmp/libomptarget/test/mapping/present/target_exit_data.c`.
> 
> I think it's wrong for `omp target data`.  My read of TR8 is that the present assertion does not happen on exit from a region.  For example:
> 
> ```
> #pragma omp target data map(alloc:x)
> { 
>   #pragma omp target data map(present,alloc:x)
>   {
>     #pragma omp target exit data map(delete:x)
>   } // fails here but shouldn't
> }
> ```
> 
In order to fix this, it looks like I could modify Clang to pass a new argument to `__tgt_target_data_end` to indicate whether it's just the end of a region (as in `omp target exit data`) or both the start and end of a region (as in `omp target data`).

Before I work on that, can somewhere confirm my understanding of TR8 behavior, described in the previous comment?


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

https://reviews.llvm.org/D83062