[Openmp-commits] [PATCH] D104418: [OpenMP][Offloading] Fixed data race in libomptarget caused by async data movement

[Guilherme Valarini via Phabricator via Openmp-commits]

gValarini added a comment.

Reading the event-based solution, one corner case came up to my mind (it is odd, but possible):

  #pragma omp target map(to: A, B) nowait
  { /* read A and B */ }
  
  #pragma omp target map(to: B, A) nowait
  { /* read A and B */ }

Note that the only difference between them (besides their internal code) is the ordering of the buffers in the mapping list.

Since both regions just read A and B, there is no need for a dependency between them (from the programmers perspective) and, thus, they can be executed by two different threads as follows:

- T1: alloc A -> issue A submission  -> create event for A -> alloc B (not new, MoveData = false)  -> wait for B event indefinitly ...
- T2: alloc B -> issue B submission  -> create event for B -> alloc A (not new, MoveData = false)  -> wait for A event indefinitly ...

This issue seems similar to the lock ordering problem. It can occur depending on the answers to the questions below:

1. Can a plugin implementation defer any command execution until the agnostic layer calls the device synchronization function (`__tgt_rtl_synchronize`)? **If yes, then the above problem can occur since no synchronization call is done.**
2. Should a plugin implementation indirectly synchronize the queue where an event was created when a call to `__tgt_rtl_wait_event` is made? **If yes, then the problem above won't happen, but this should be clarified in the documentation of the API.**
3. Are the variables in the mapping list ordered somehow? **If that is the case, then the (possible) problem described above will never happen as well (like resolving it through lock ordering).**

If this is a problem, sorting the mapping list or ensuring the behavior described in the second question would probably solve it. What do you think? Is this line of thought correct?


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D104418