In D85692#2228860, @nickdesaulniers wrote:

In D85692#2227531, @jdenny wrote:

Does this mean we officially no longer support python 2, which this change breaks?

Send a patch to support both and I'll approve it.

Does this mean we officially no longer support python 2, which this change breaks?

Thanks!

Don't examine LIBOMPTARGET_DEBUG=1 output in the test. It doesn't work in release builds.

Thanks for the quick reviews!

Add user diagnostic for mapping extension. Use PRId64 where appropriate.

This patch fixes DeviceTy::getTgtPtrBegin to return null when
target_data_update is the caller. I'm wondering if it should do the
same for more callers.

In D85246#2194893, @jdoerfert wrote:

Maybe change the commit title to OpenMP to confuse less people.

Thanks for the quick review.

Thanks for the review.

Thanks for the review.

Adjust the fixme as requested.

Thanks!

Should this patch be marked NFC?

Adjusted logic for rejecting present as requested. Rebased.

Thanks!

Replaced SeparateBeginEnd parameter with new TargetDataInfo field as requested. Rebased.

Are you following an existing TableGen emitter as a precedent for this kind of thin wrapper design?

Sure. Thanks for the reply.

Thanks!

In D84557#2178019, @protze.joachim wrote:

Broken because of D83963, @jdenny proposed a fix, which works for me:

libomptarget :: mapping/alloc_fail.c
libomptarget :: mapping/present/target.c
libomptarget :: mapping/present/target_data.c
libomptarget :: mapping/present/target_enter_data.c
libomptarget :: mapping/present/target_exit_data.c
libomptarget :: mapping/present/zero_length_array_section.c
libomptarget :: mapping/present/zero_length_array_section_exit.c

I've pushed it to master as 9b4826d18b5f.

not expects a non-zero exit not a signal. Adding --crash adjusts the expectation.

Ping.

Rewrite patch as discussed: instead of generating different runtime calls for the end of an omp target data vs. the beginning of an omp target exit data so that the runtime can determine when to ignore present, change Clang to filter present from the map types at the end of an omp target data.

In D84422#2176802, @grokos wrote:

In D84422#2173500, @jdenny wrote:

I've added a comment to the runtime code that performs the check. As you can see, the check is performed regardless. It's just a question of whether the runtime treats it as an error. I don't think performance is an issue.

My concern here is that it will be hard to justify changes to the runtime if I cannot formulate a use case.

Thinking about it, I don't think there can be a case where something is present upon entering a target region and not be present when we're exiting. Whatever code comprises the target region is code executed on the device - it cannot modify the state of host objects (i.e. libomptarget) in any possible way. E.g. the kernel cannot invoke libomptarget functions, allocate memory, map/unmap data etc.

The only case where something like this would be possible is if we have multiple host threads executing async offloading. In such a case, one thread may launch a target region at a moment when the requested mapping is present on the device and while the kernel is executing some other thread performs a target data exit on the desired mapping. Upon exiting the kernel, the mapping will no longer be present but this is clearly a race condition (user's fault), so I don't think we should pay attention to such a scenario.

Other than the nit I just added about comment duplication, this LGTM. Thanks. However, let's a wait a couple of days in case anyone else has an opinion.

This change broke a number of libomptarget tests. Adding --crash to the %libomptarget-run-fail-* definitions in openmp/libomptarget/test/lit.cfg fixes that for me. Any objections?

In D84422#2173449, @RaviNarayanaswamy wrote:

In D84422#2173372, @jdenny wrote:

In D84422#2172898, @jdenny wrote:

Has anyone clarified the motivation for this behavior?

I meant, is there any insight into why the spec specifies this behavior?

In D84422#2172926, @grokos wrote:

Instead of introducing new API functions and making all these changes in all these files, wouldn't it be easier if we just unset the PRESENT flag from arg_types in clang when we generate the call to __tgt_target_data_end_* if we are exiting from a scoped environment?

Ah, that does sound simpler. Thanks. I'll look into it.

Suppressing the presence check on exit from omp target would require a runtime change in addition to the Clang change you suggest for omp target data. However, I've so far failed to formulate a reasonable test case. Specifically, I don't yet see a way to guarantee that the data will definitely be present at the start of omp target but might not be present by the end. Is it possible? If not, then maybe we should leave the check in place for omp target.

I would rather not have a check if not required by the spec as it would an unnecessary overhead to performance.

In D84422#2172898, @jdenny wrote:

Has anyone clarified the motivation for this behavior?

In D84422#2170702, @RaviNarayanaswamy wrote:

In D84422#2170667, @grokos wrote:

So is the test case that motivated this patch illegal OpenMP code?

#pragma omp target enter data map(alloc:i)
#pragma omp target data map(present, alloc: i)
{
  #pragma omp target exit data map(delete:i) // you cannot delete that object in the scope, illegal code?
} // fails presence check here

According to spec the test should work. ie should not check for presence on exit from a blocked openmp pragma scope.

Thanks for working on this. I've been wanting to see this fixed too. I agree that lit is the culprit.

LGTM

I don't know if the OpenMP committee has any documented rationale for this behavior. I can say that the OpenACC committee is considering the same semantics. However, the issues to consider are not identical. For example, OpenACC has a separate structured reference counter, meaning it should be impossible for such data not to be present at the exit of a data construct unless you've shut down the runtime.

Rebased. Due to D84182, updated the map types expected by my new tests in clang/test/OpenMP/target_data_codegen.cpp.

Thanks again for the review.

Tried to fix comments on PTR_AND_OBJ.