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

[Alexandre Eichenberger via Phabricator via Openmp-commits]

AlexEichenberger added a comment.



> In https://reviews.llvm.org/D50522#1194903, @Hahnfeld wrote:
> 
>> Suppose we have 2 devices plugged into the system, and the first one cannot be used (for whatever reason: hardware failure, exclusive configuration and somebody else is running, etc.).
>>  Now a clever application sees the two (because `omp_get_num_devices()` returns 2) and does:
>>
>>   #pragma omp parallel num_threads(omp_get_num_devices())
>>   {
>>     #pragma omp target device(omp_get_thread_num())
>>     { }
>>   }
>>
>>
>> I think the runtime behaviour with this patch depends on the execution order (and exposes a race condition in `handle_target_outcome` on `TargetOffloadPolicy`; let's ignore that for now):
>>
>> - If `target device(0)` executes first, libomptarget will notice the error and silently disable offloading. All `target` regions will execute on the host.
>> - If however `target device(1)` executes first and returns successfully, libomptarget will raise `OMP_TARGET_OFFLOAD` to `MANDATORY` and will abort execution when catching the error of `target device(0)`.
>>
>>   I don't think that makes much sense. IMO the runtime should detect two "visible" -> "available" devices and abort execution in all cases.
> 
> 
> Did you consider this example?

Your example can show a different issue as well:

A program first do a register_lib that has code for a device which does not exist on that machine. Thus omp_get_num_devices is zero. DEFAULT then becomes DISABLED. A program then does a register_lib that has code for a device that exist on this machine. But now it is disabled.

A way out of this is to see the bigger picture. If the user want to guarantee execution of all targets on a device, the user must us MANDATORY. If the user does not want devices, then DISABLED is called for. DEFAULT is a best effort, it will not be perfect, nor it has to be perfect.

I suggest that deciding on the first attempt to actually offload is as valid a policy as any, and is simple to understand. Regardless of how we do it, what we want to really avoid is that we execute some kernels on the device, and some on the host, (ignoring here explicit orders from the program via the "if(0)" clause). Both policies "num_devices>0" and "decide on first invocation" satisfy this.


Repository:
  rOMP OpenMP

https://reviews.llvm.org/D50522