LGTM

LGTM

LGTM

fantastic

response to Alexey's comment

LGTM, just made one minor suggestion

see note above; apologies if it is already done and hiding somewhere I did not see

See suggested changes, should be pretty straightforward to do, let me know if you need help

see comments

make sense, thanks

Yes, we should have done this right away. I did a rewrite of the functionality that uses header files already included; it was posted in D50522 and the changes were approved and pushed to svn.

• Merge branch 'unpatched-master' into unpatched-master-target-offload-envvar
• while (0) edit

• added do while(1) for macros

• put fatal message for default tgt test

Alexey was not happy about adding an include, and we already used the define approach for the debug messages. So we reuse the same functionality that we already used for debug.
Agreed that Johnas suggested it, and given the problem with the include, Alexey urged me to change the code to what we have here.

In D51226#1212685, @Hahnfeld wrote:

I don't understand this: In the last review you said that we need the complicated FatalMessage and now you are reverting that?!?

• format

I reworked the solution to avoid varargs, and proposed a more robust solution in D51226

In D51107#1211035, @Hahnfeld wrote:

Thanks to all for your valuable comments, much appreciated, really contributed to the quality of the patch

updated comments status

fixed clang-format and moved fatal message function

responded to comments

implemented suggest changes

responded to comments.

Addressed all remaining issues

Good arguments, I see why the consistency of why linking the number of devices returned and OMP_TARGET_OFFLOAD environment is good.
Implemented the requested changes

Default policy is now selected when registering the libraries, and will default to mandatory/disabled depending of whether there are one or more devices/none.

Added a mutex around changing DEFAULT to MANDATORY or DISABLED
Added proper action when discovering failure (which is immediate abort of function being performed)
Failure to locate data during an update is now tolerated (a warning could be issued)

In 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?

I disagree. As it is currently written, omp_get_num_devices() can also grow over time, so you can also have it return zero, only to increase later to a larger number. What does that do to DEFAULT?
I believe the current policy is ok; if any device fails for any reason on first use, it becomes disabled. Anyone that want to rely on devices being there should use the MANDATORY policy.
I am happy to add a lock around the change of DEFAULT to MANDATORY or DISABLED.

In D50522#1194210, @Hahnfeld wrote:

You are right, "available" is not defined in the standard. I've always though of "plugged into the system", ie all devices that are visible to the CUDA runtime. That would match the current implementation of omp_get_num_devices which is defined to return "the number of available target devices".
Actually this behaviour would be important for us as we have our GPUs configured exclusively. So when there is already a process running all other users get a runtime error. In that case it would be very helpful to have libomptarget abort the program.

In D50522#1194100, @Hahnfeld wrote:

Does this patch supersede D44522?

agreed

First, there missing sync for the _tgt_target_data_begin_depend, tgt_target_data_begin_depend_nowait, tgt_target_data_update_depend, tgt_target_data_update_depend_nowait,
Look at the end versions of the call for a possible implementation.
Second, it would be better IMO to use kmpc_omp_wait_deps instead of the broader task wait.

Changes are correct, a memory barrier is needed between allocating/reseting values to zero and publishing the pointer. That way, another thread see the published pointer only once all of the zeroing is complete.