In D87841#2280009, @jdoerfert wrote:

In D87841#2279743, @lebedev.ri wrote:

As openmp and llvm are now developed in the same monorepo, the only uses this
dependency can break are those that download the monorepo, then delete the llvm
directory before calling cmake. That is not a compelling use case.

That is factually not accurate.
There are standalone tarballs, including one for omp:
https://releases.llvm.org/download.html#10.0.1 OpenMP Source code (.sig)

Right. You can download the source for the sub-projects.
Though, others, e.g., clang, are already not compileable in isolation.

I see three options here:

1. Continue to copy code from llvm-core to openmp and back. This is "good" for people that compile libomp in isolation and "bad" for developers. Arguably, it is also "bad" for people that use libomp with clang because there is an increased risk we have an unnoticed mismatch between the two. This is not theoretical but already happened.
2. Copy the code from llvm-core to openmp during packaging time. This is a (small?) burden on the package maintainers but has all the advantages for now. However, it will get more complicated the more llvm-core features we want to use, e.g., tablegen.
3. Require users to download llvm-core and point to it. This is a small burden on the users but has all the advantages now and in the future.

This patch, on its own, picks 3). I'm happy with that.

Minor comments (about typos) to be taken into account in case of a future patch.

Looks good!

Nice catch! Looks good.

LGTM. Thanks!

Looks good.

OK, now it makes sense. LGTM

I suppose the same must be applied to all cases where present is used. Without a modification like this, present just confirms that the starting address of the object is already mapped, although a size may be specified in the map clause.

#pragma omp target map(present: p[0:10]) // present will succeed even if only p[0] is mapped and p[1]-p[9] are not

So the question here is: should present apply to the size as well as the begin address? I would say yes.

So does the mapper function emit entries in reverse order upon exiting a target/target data region?

In D68100#2173822, @lildmh wrote:

This is an optimization brought up by Deepak. I guess you were in that meeting too but forgot. It could be quite useful when you map an array of struct/class. Assume you map 1000 of this structure, with this optimization most memory allocation can be done in a single allocation, instead of allocation 12 bytes memory 1000 times.

Thinking about it, it's actually important for correctness too. Assume you map C a[2]. If you map separately, a[0] and a[1] could be mapped to not contiguous locations, and it will cause error/segfault when GPU kernel access this array. If you allocate the whole array a[2] together, such problem won't happen.

Thanks, it makes sense. Can you submit a patch?

Any idea why this change broke some tests? And how does --crash fix the problem?

This looks much better now. I don't have any other comments. Since this patch is now essentially a clang-only patch, I'll let @ABataev accept it or post comments.

(Earlier I posted here an answer for another patch). Looks good, obviously it must be committed after the clang patch.

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.

In D68100#2173612, @lildmh wrote:

The first combined entry comes from mapping the whole structure. I think because of the alignment, the structure is actually 16 bytes. The 2nd combined entry is the real entry emitted to map the structure. Why it looks like there are 2 of them, because at the beginning of a mapper function, it maps the whole structure no matter what, which generate the 1st combined entry you saw here. Then we generate detailed mapping entry, which generates the 2nd combined entry you saw here. They are not necessarily the same. It happens to be similar in this example.

After this patch was committed, I tried to run the following example:

#include <stdio.h>

In D68100#2173460, @lildmh wrote:

Thanks George for looking into this, and sorry for the late response.

I believe this is not a bug, it's a design choice we made early. The design choice is we map the whole structure at the beginning for one piece so we don't map the individual parts of them separately, which may cause a lot of memcpy.

For the RefCount, when the runtime check the 2nd component in your example, it will find it's already mapped and will not increase the RefCount, so I think it's not a bug and the behavior is correct.

So let's proceed with the patch.

@lildmh: I think I've found a bug. I used declare_mapper_target.cpp. When we call the mapper function, it generates 3 components. The first 2 are identical and correspond to the parent struct. This is what MapperComponents looks like inside function target_data_mapper:

(gdb) print MapperComponents
$1 = {Components = std::vector of length 3, capacity 4 = {{Base = 0x7fffffffb598, Begin = 0x7fffffffb598, Size = 8, Type = 32}, {Base = 0x7fffffffb598, Begin = 0x7fffffffb598, Size = 8, Type = 32}, {Base = 0x7fffffffb598,
      Begin = 0x62efd0, Size = 4096, Type = 562949953421331}}}

Mapping the parent struct twice is problematic. If we have more struct members and some of them are NOT pointers, then upon target_data_end libomptarget will check the parent struct's reference counter to determine whether the scalar member must be copied back to the host. If the reference counter is greater than 1, then the runtime will skip copying back the scalar. Mapping the parent struct two times in a row results in RefCount=2.

I tried the patch and indeed it fixes the problem with target_depend_nowait.cpp. I'll let someone else do the review though.

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

What confuses me about this interpretation of the standard is the inconsistency at data exit. So if we have an explicit omp target exit data map(present...) then we should respect the "present" semantics, whereas when we have a scoped data exit:

#pragma omp target data map(present,...)
{
  ...
} // implicit "exit data" here

then "present" should be ignored.

@lildmh I've got a question unrelated to the problem we are discussing here. I ran declare_mapper_target.cpp and when libomptarget calls the mapper function it prints the following:

Libomptarget --> __tgt_push_mapper_component(Handle=0x00007ffcebd0fb48) adds an entry (Base=0x00007ffcebd101e0, Begin=0x00007ffcebd101e0, Size=8, Type=0x20).
Libomptarget --> __tgt_push_mapper_component(Handle=0x00007ffcebd0fb48) adds an entry (Base=0x00007ffcebd101e0, Begin=0x000000000231bfe0, Size=4096, Type=0x2000000000013)

Why is the second entry's MEMBER_OF field set to 2? It should be MEMBER_OF 1, since the pointer-pointee pair c.a[0:N] is part of struct c which is the first entry on the list.

OK, I suspect there is a race condition involving the CUDA plugin. If I compile the test on x86_64-pc-linux-gnu then I always get the correct result, no matter whether we print debug output or not.

My system was messed up and used libraries and compilers from different builds. Please ignore my previous message.

I tried to run declare_mapper_target.cpp on a Nvidia GPU. The problem occurs while loading the device image:

Target CUDA RTL --> Error returned from cuModuleLoadDataEx
Target CUDA RTL --> CUDA error is: device kernel image is invalid

Sounds like a clang problem, I don't see why libomptarget could be the culprit here.

Sure, go ahead!

OK, let's proceed with this version now and later on we can fine-tune the behavior.

I'm having a look. The problem actually occurs after D67833 (the clang patch), not D68100. If you see at the backtrace in gdb, control never reaches a libomptarget entry point, things fail once we enter the outlined function for the parallel region via __kmp_invoke_microtask. I doubt it's a libomptarget issue.

abort() does not flush stdout, but here we are printing to stderr which is not buffered so there is no need to flush.

I can confirm the warnings are gone with this fix.

OK, now it works. Thanks!

I tried to build clang with this patch and I get errors like:

CGOpenMPRuntime.cpp:9463:38: error: ‘OMPRTL___tgt_target_teams_nowait_mapper’ was not declared in this scope
                                    ? OMPRTL___tgt_target_teams_nowait_mapper

No, it's not needed anymore. This patch bypasses the need to do that refactoring. Can you please abandon that revision?

Ping. If the patch lands toady or tomorrow, then we will meet the clang-11 deadline and include support for declare mapper.

The patch looks good now, thanks for all the work!

I tried to address our previous complaints about code duplication and came up with a scheme which results in a much shorter and cleaner diff with virtually no code duplication. Instead of refactoring code from taget_data_begin/end/update, I introduced a new internal function target_data_mapper which generates new arrays args_base, args, arg_sizes and arg_types for the custom mapper and calls target_data_begin/end/update again using the new arguments.

Looks good. Do we know why the test doesn't run without -fPIC?

Like D82224, looks good.

In D82718#2124420, @tianshilei1992 wrote:

I did some investigation and finally think BLOCKING_SYNC might be a good option here, but I also would like to hear from others.
Basically we have three options here: SPIN, YIELD and the BLOCKING_SYNC.

• SPIN: In most cases this is not a good option but it might be better for a really tiny kernel.
• YIELD: Like mentioned in CUDA documentation, it can increase latency, but can increase the performance of CPU threads performing work in parallel with the GPU. But for OpenMP, this level of yield may not help too much because it is not very common to have thread oversubscription in OpenMP.
• BLOCKING_SYNC: I guess chances are that it works with interruption. This is kind of a balanced option which will not increase too much latency, and also will not waste the resource, especially considering that we might be going to use unshackled threads for all target tasks.
• AUTO: In fact this option will be used if there is no flag specified. According to CUDA documentation, the behavior depends on the number of CUDA contexts in current process versus the number of logical processors in the system. Either SPIN or YIELD will be used. BLOCKING_SYNC will only be chosen on Tegra devices.

Looks good now. Thanks!

The patch looks good now. I don't have any other comments. We can revisit the driver flag proposal for setting the file extension in the future.

I'm happy with the changes, I think the patch looks good now.

The partial linking scheme has been found to not work correctly in all cases (it fails when we have libraries with device code only). A new patch will be uploaded which will be based on archive extraction.

LGTM.

I think it would look better if we initialized those variables inside Device::getOrAllocTgtPtr - the caller should not assume any default value. That's also what happens in target_data_end with the IsLast variable, it is initialized inside Device::getTgtPtrBegin.

I guess it's a reasonable thought to disable such verbose output by default, but I can also see why these diagnostics could be useful l in some cases. Can we introduce a CMake variable to toggle --ptxas-options=-v on and off? (off by default)

Superseded by D74145. You can abandon this one.

Nice catch! Thanks!

Done!

Building the CUDA plugin on aarch64 hosts was enabled in the ykt branch ages ago, so the configuration has been tested extensively. LGTM.

I, too, can confirm that this patch fixes the problem. It seems we need to think more about how libomptarget is initialized to avoid similar problems in the future.

@ABataev, @jdoerfert, does the patch look good to you? Can someone accept it if it's ready to go?

Looks good. Do you think I should also change the priorities in my clang-offload-wrapper patch to be out of the reserved range?
https://reviews.llvm.org/D75223

I think we should move forward (I still haven't been able to reproduce the problem). LGTM.

In D74837#1883479, @ABataev wrote:

Maybe it is just the problem of the test. I have CHECK-NOT in it to detect a possible crash. Maybe, just missed a variant for your system. Did you try the reproducer from the error report?