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

[Joel E. Denny via Phabricator via Openmp-commits]

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:
+      //
----------------
jdenny wrote:
> grokos wrote:
> > jdenny wrote:
> > > 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.
> > Yes, we can include an assertion here.
> The assert would currently fail for the following:
> 
> ```
> int *x;
> #pragma omp declare target link(x)
> 
> int main() {
>   #pragma omp target map(present, tofrom: x[0:3])
>   ;
>   return 0;
> }
> ```
> 
> 
I didn't add the assert.  I updated the comments to describe the situation as I understand it.  Please advise if it's not right.


================
Comment at: openmp/libomptarget/src/omptarget.cpp:416
+                DPxPTR(HstPtrBegin), data_size);
+        return OFFLOAD_FAIL;
+      }
----------------
jdenny wrote:
> 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?
> 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).

I reversed "omp target exit data" and "omp target data" there.  Otherwise, I believe that comment is correct.

For now, I've decided to add a fixme and handle this in a later patch after pushing this series.


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

https://reviews.llvm.org/D83062