[Openmp-commits] [PATCH] D50522: [OpenMP][libomptarget] Bringing up to spec with respect to OMP_TARGET_OFFLOAD env var

[Joachim Protze via Phabricator via Openmp-commits]

protze.joachim added a comment.

@Hahnfeld and I discussed the behavior and found a very inconsistent behavior. Let's consider something like under the constraint, that device(1) is busy (will fail to offload) and device(0) is free:

In https://reviews.llvm.org/D50522#1194903, @Hahnfeld wrote:

>   #pragma omp parallel num_threads(omp_get_num_devices())
>   {
>     #pragma omp target device(omp_get_thread_num())
>     { }
>   }
>


We found the following cases:

1. first offload to device(0) -> succeed -> set MANDATORY -> abort on offloading to device(1)
2. first offload to device(1) -> fail -> set DISABLED -> continue execution on the host
3. start offloading to device(0) and device(1) -> fail on 1 -> set DISABLED, but execution of offloaded code will complete. -> although we are in DISABLED, we executed code on a device

The locking only prevents the race on the variable, but not the behavior in 3).

>From our point of view, the decision between DISABLED and MANDATORY must be before starting any offloading.

Another reason to implement such behavior:
If the offloaded code changes the memory state (and the target region code is not idempotent), this prevents inconsistent states of execution, if the initial offloading fails after partial execution on the device and a second execution the host.


Repository:
  rOMP OpenMP

https://reviews.llvm.org/D50522