I'll add a test to it. Please give me some time as I am busy with other stuffs.

Your understanding above is exactly right. It should be equivalent to

#pragma omp target data map(tofrom: c)
#pragma omp target data map(tofrom: c.a[0:NUM])
#pragma omp target teams distribute parallel for
for (int i = 0; i < NUM; i++) {
  ++c.a[i];
}

In D68100#2173949, @protze.joachim wrote:

From my perspective, the declare_mapper_target.cpp code is semantically equivalent to:

#pragma omp target data map(tofrom: c)
#pragma omp target data map(tofrom: c.a[0:NUM])
#pragma omp target teams distribute parallel for
for (int i = 0; i < NUM; i++) {
  ++c.a[i];
}

and

#pragma omp target enter data map(to: c)
#pragma omp target enter data map(to: c.a[0:NUM])

#pragma omp target teams distribute parallel for
for (int i = 0; i < NUM; i++) {
  ++c.a[i];
}
#pragma omp target exit data map(from: c.a[0:NUM])
#pragma omp target exit data map(from: c)

Can you express the behavior of your mapping implementation in terms of OpenMP target enter/exit data primitives?

In D68100#2173703, @grokos wrote:

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.

I assure you that's not how structs are mapped.

You don't map "the whole struct", you only map what is needed. For this you emit a combined entry which has a size large enough to encompass all members we are interested in. Then entries for individual members follow. One combined entry + as many member entries as needed. The first entry which "maps the whole struct" should not be there and is plain wrong.

In D68100#2173541, @grokos wrote:

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.

No, this is not related to our design choices. Here we are mapping the whole struct twice for no reason. The entries should be:

1) combined entry (i.e. the entry that maps the whole struct
    base = &c, begin = &c.a, size = sizeof(class S), type = TARGET_PARAM
2) member entry for c.a[0:NUM]
    base = &c.a, begin = &c.a[0], size = NUM*sizeof(int), type = MEMBER_OF(1) | PTR_AND_OBJ | TO | FROM
3) member entry for c.onHost
    base = &c, begin = &c.onHost, size = sizeof(int), type = MEMBER_OF(1) | TO | FROM

But what happens now is that the combined entry is emitted twice, so MapperComponents looks like this:

<combined entry>, <combined entry>, <entry for c.a[0:NUM]>, <entry for c.onHost>

instead of

<combined entry>, <entry for c.a[0:NUM]>, <entry for c.onHost>

And what's more, the first combined entry has size=16 whereas the second combined entry has size=12. Where does this 16 come from? The size of the struct is 12 bytes (a pointer + an int). This also explains why the MEMBER_OF field is set to 2, because the second element in the list of arguments is also the combined entry.

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

The problem you had is probably not related to this patch. A lot of libomptarget tests are not working for nvptx now. I suppose someone is going to fix them.

Can you check if https://reviews.llvm.org/D84470 fixes the problem as this test doesn't work on my machine? If yes, please accept and upstream it since I don't have the permission.

In D68100#2168350, @grokos wrote:

@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.

I couldn't reproduce this problem because all commands are ignored on my machine. Could you send the exact compiling and running commands?
In the IR files you sent me before, any idea where segfault happens?

Thanks. Another weird place is it passes with nvprof. Not sure why using nvprof makes a difference here.

I can reproduce this. When running the test itself, Sum=1024.
When running the test with nvprof, Sum=2048. Combining with that you said Sum=2048 when LIBOMPTARGET_DEBUG=1, I suspect the GPU offloading is disabled in the above case. Any idea what happened recently to libomptarget which can potentially cause this problem? I didn't follow the recent development so have no idea.

In D68100#2167208, @protze.joachim wrote:

The other tests failed after the commit. They started to succeed with various commits.
At above mentioned commit, only this single test fails.

In D68100#2166450, @protze.joachim wrote:

declare_mapper_target.cpp still fails for me consistently in RUN line 2, so for the nvptx version. I execute on x86_64 with Tesla V100 and cuda 10.0.
When I execute the test with export LIBOMPTARGET_DEBUG=1, the test succeeds.
In case of failure, the test prints Sum = 1024, in case of success, the test prints Sum = 2048 as expected.

I run the tests with a fresh build (ae31d7838c36).

In D67833#2166478, @protze.joachim wrote:

Starting with the commit of this patch, the libomptarget test llvm-project/openmp/libomptarget/test/offloading/target_depend_nowait.cpp fails. Here is a stacktrace of the segfault:

$ gdb target_depend_nowait.cpp.tmp-x86_64-pc-linux-gnu
(gdb) run
...
Program received signal SIGSEGV, Segmentation fault.
(gdb) bt
#0  0x0000000000400fcf in .omp_outlined._debug__ (.global_tid.=0x2aaab96b9be0, .bound_tid.=0x2aaab96b9bd8) at target_depend_nowait.cpp:22
#1  0x0000000000401e8d in .omp_outlined..23 (.global_tid.=0x2aaab96b9be0, .bound_tid.=0x2aaab96b9bd8) at target_depend_nowait.cpp:18
#2  0x00002aaaab574213 in __kmp_invoke_microtask () from libomp.so
#3  0x00002aaaab51338e in __kmp_invoke_task_func () from libomp.so
#4  0x00002aaaab5126bf in __kmp_launch_thread () from libomp.so
#5  0x00002aaaab55d3d0 in __kmp_launch_worker(void*) () from libomp.so
#6  0x00002aaaabbd6e65 in start_thread () from libpthread.so.0
#7  0x00002aaaabee988d in clone () from libc.so.6

In D67833#2161241, @Meinersbur wrote:

Seems that it already has been applied ;-)

Thanks George!

Fix is at https://reviews.llvm.org/D83959
@Meinersbur Michael, could you please help upstream the patch above? Thanks

In D67833#2155196, @fhahn wrote:

I think this change introduced the following warnings, where auto & is used for types that are always copied. It would be great if you could take a look.

lvm-project/clang/lib/CodeGen/CGOpenMPRuntime.cpp:8011:24: warning: loop variable 'L' is always a copy because the range of type 'clang::OMPMappableExprListClause<clang::OMPToClause>::const_component_lists_range' (aka 'iterator_range<clang::OMPMappableExprListClause<clang::OMPToClause>::const_component_lists_iterator>') does not return a reference [-Wrange-loop-analysis]
      for (const auto &L : C->component_lists()) {
                       ^
llvm-project/clang/lib/CodeGen/CGOpenMPRuntime.cpp:8011:12: note: use non-reference type 'std::__1::tuple<const clang::ValueDecl *, llvm::ArrayRef<clang::OMPClauseMappableExprCommon::MappableComponent>, const clang::ValueDecl *>'
      for (const auto &L : C->component_lists()) {
           ^~~~~~~~~~~~~~~
llvm-project/clang/lib/CodeGen/CGOpenMPRuntime.cpp:8016:24: warning: loop variable 'L' is always a copy because the range of type 'clang::OMPMappableExprListClause<clang::OMPFromClause>::const_component_lists_range' (aka 'iterator_range<clang::OMPMappableExprListClause<clang::OMPFromClause>::const_component_lists_iterator>') does not return a reference [-Wrange-loop-analysis]
      for (const auto &L : C->component_lists()) {
                       ^
llvm-project/clang/lib/CodeGen/CGOpenMPRuntime.cpp:8016:12: note: use non-reference type 'std::__1::tuple<const clang::ValueDecl *, llvm::ArrayRef<clang::OMPClauseMappableExprCommon::MappableComponent>, const clang::ValueDecl *>'
      for (const auto &L : C->component_lists()) {
           ^~~~~~~~~~~~~~~
llvm-project/clang/lib/CodeGen/CGOpenMPRuntime.cpp:8032:24: warning: loop variable 'L' is always a copy because the range of type 'clang::OMPMappableExprListClause<clang::OMPUseDevicePtrClause>::const_component_lists_range' (aka 'iterator_range<clang::OMPMappableExprListClause<clang::OMPUseDevicePtrClause>::const_component_lists_iterator>') does not return a reference [-Wrange-loop-analysis]
      for (const auto &L : C->component_lists()) {
                       ^
llvm-project/clang/lib/CodeGen/CGOpenMPRuntime.cpp:8032:12: note: use non-reference type 'std::__1::tuple<const clang::ValueDecl *, llvm::ArrayRef<clang::OMPClauseMappableExprCommon::MappableComponent>, const clang::ValueDecl *>'
      for (const auto &L : C->component_lists()) {
           ^~~~~~~~~~~~~~~
3 warnings generated.

In D67833#2154312, @grokos wrote:

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

Where are these OMPRTL___tgt_ symbols defined?

Include the llvm part

@grokos Could you upstream this patch and the runtime patch neck to neck? Upstreaming one of them will break the OpenMP offloading. It will be nice if you can test this patch locally with the runtime patch. Thanks

Update diff and pass test. Please accept

I'll update the diff and please check and accept after that

In D68100#2151386, @grokos wrote:

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

Also, not sure if https://reviews.llvm.org/D71782 is still needed. Please check

The codegen patch needs to land before this one to pass these test programs

Test cases will be uploaded in another patch when the Clang patch is upstreamed. That Clang patch depends on this (https://reviews.llvm.org/D67833). So I think the order is this patch, clang patch, test patch.

Address Jon's comments

Rebase and rediff with the nfc version

Yes, this is the NFC part of the declare mapper patch. I've tested it. There is not much test in libomptarget though.

Change the function name

Address Alexey's comments to change mapper function size and refactor code

Alexey, thanks for the review

I'll split the patch into 2 later

Rebase and address comments

Thanks for your reviews. Hope this looks better.

Rebase

Thanks Alexey and Jon for your review. Fixed the issues and rebased

In D68100#1686458, @grokos wrote:

Lingda is right, we had faced the same issue in the loop trip count implementation. The loop trip count should be set per task but libomptarget has no notion of tasks, so we ended up engaging the host runtime (libomp) to store per-task information. Although it involves more work, I still believe that will be a more elegant solution.

Alexey and George: This is a big decision to make. We need to have most people's consents. I'll send it to the mailing list later.

In D68100#1686437, @ABataev wrote:

In D68100#1686432, @lildmh wrote:

Sorry, you are right. I didn't think about the case to always clean up the mappers after finishing using it.

Another possible problem: what if a task is scheduled out after __tgt_mapper and before __tgt_target, for example. I don't think we can keep a per task/thread mapper storage in the current implementation.

tgt_mapper must be called immediately before tgt_target in the same task context.

Sorry, you are right. I didn't think about the case to always clean up the mappers after finishing using it.

In D68100#1686413, @lildmh wrote:

In D68100#1686397, @ABataev wrote:

In D68100#1686395, @lildmh wrote:

In D68100#1686337, @grokos wrote:

In D68100#1686044, @ABataev wrote:

In D68100#1685975, @lildmh wrote:

Please review when you have time, thanks

Maybe, it is better to add a single function to pass the list of mappers to the runtime and use the old functions rather than just duplicate them? Something like __tgt_mappers(...) and store them in the runtime. Plus, modify the original functions to use mappers if they were provided. Thoughts?

I'm in favor of this solution, as it is less intrusive than the one posted in this patch. I think a revised patch will be quite smaller and the changes it introduces will be clearer.

I think Alexey and George's concern is that this patch introduces many new runtime apis. My arguments are:

1. This patch will deprecate the old interfaces as we discussed before in the meeting. E.g., __tgt_target_teams will no longer be used. They are kept there just for legacy code support. So I didn't actually introduce any new runtime apis.

2. If we have something like __tgt_mapper instead, and integrate all mapper functionalities in it, we will need to pass extra arguments to it to distinguish whether it is for __tgt_targt, __tgt_target_data_begin, etc. On the other hand, the current runtime has an interface for each case. For instance, we have __tgt_targt, __tgt_target_data_begin, __tgt_target_data_update_nowait, etc., instead of a single __tgt_target function which does it all. Therefore, I think the current design fits the overall design of OpenMP runtime better: have a function for each case.

Why we will need this extra argument? It just provides a list of mapper functions and stores them in the runtime before we call any __tgt_... function. Each particular __tgt_... runtime function will know what do with all these mappers if they were stored.

Okay, I think I understand your idea now. Then in this case, we will have a call to __tgt_mapper before every call to __tgt_target*, because we need to overwrite the mappers written for previous calls. I don't particularly like this idea, since this will introduce implicit dependencies between different runtime calls and a program will have twice runtime calls. But if most people like it, I'm okay.

In D68100#1686397, @ABataev wrote:

In D68100#1686395, @lildmh wrote:

In D68100#1686337, @grokos wrote:

In D68100#1686044, @ABataev wrote:

In D68100#1685975, @lildmh wrote:

Please review when you have time, thanks

Maybe, it is better to add a single function to pass the list of mappers to the runtime and use the old functions rather than just duplicate them? Something like __tgt_mappers(...) and store them in the runtime. Plus, modify the original functions to use mappers if they were provided. Thoughts?

I'm in favor of this solution, as it is less intrusive than the one posted in this patch. I think a revised patch will be quite smaller and the changes it introduces will be clearer.

I think Alexey and George's concern is that this patch introduces many new runtime apis. My arguments are:

1. This patch will deprecate the old interfaces as we discussed before in the meeting. E.g., __tgt_target_teams will no longer be used. They are kept there just for legacy code support. So I didn't actually introduce any new runtime apis.

2. If we have something like __tgt_mapper instead, and integrate all mapper functionalities in it, we will need to pass extra arguments to it to distinguish whether it is for __tgt_targt, __tgt_target_data_begin, etc. On the other hand, the current runtime has an interface for each case. For instance, we have __tgt_targt, __tgt_target_data_begin, __tgt_target_data_update_nowait, etc., instead of a single __tgt_target function which does it all. Therefore, I think the current design fits the overall design of OpenMP runtime better: have a function for each case.

Why we will need this extra argument? It just provides a list of mapper functions and stores them in the runtime before we call any __tgt_... function. Each particular __tgt_... runtime function will know what do with all these mappers if they were stored.

In D68100#1686337, @grokos wrote:

In D68100#1686044, @ABataev wrote:

In D68100#1685975, @lildmh wrote:

Please review when you have time, thanks

Maybe, it is better to add a single function to pass the list of mappers to the runtime and use the old functions rather than just duplicate them? Something like __tgt_mappers(...) and store them in the runtime. Plus, modify the original functions to use mappers if they were provided. Thoughts?

I'm in favor of this solution, as it is less intrusive than the one posted in this patch. I think a revised patch will be quite smaller and the changes it introduces will be clearer.

Please review when you have time, thanks

Rebase

Thanks Alexey! Please check the other 2 mapper patches at https://reviews.llvm.org/D67833 and https://reviews.llvm.org/D68100 when you have time. They should be the last mapper patches.

Fix mapper type checking

In D67978#1684169, @ABataev wrote:

In D67978#1684104, @lildmh wrote:

In D67978#1683166, @ABataev wrote:

In D67978#1683146, @lildmh wrote:

HI Alexey, the ast print test is already there. Because I didn't check the mapper for array type before, such code will always not report any error, and ast print test is correct. Codegen test belongs to the other patch I released. It fits that patch much better.

How is this possible? If we did not have support for the array type, we could not have correct handling of such types in successful tests.

The ast print for array with mapper was correct because the mapper id is still with the array type. Without this patch, the problem is it will not look up the mapper declaration associated with the id, and as a result, the codegen is not correct. I found this problem when I tested the codegen.

Then another one question. Why we don't emit the diagnostics if the original type is not a class, struct or union? We just ignore this situation currently but we should not. The error message must be emitted. I think that's why the ast-print test works correctly though it should not.

In D67978#1683166, @ABataev wrote:

In D67978#1683146, @lildmh wrote:

HI Alexey, the ast print test is already there. Because I didn't check the mapper for array type before, such code will always not report any error, and ast print test is correct. Codegen test belongs to the other patch I released. It fits that patch much better.

How is this possible? If we did not have support for the array type, we could not have correct handling of such types in successful tests.

HI Alexey, the ast print test is already there. Because I didn't check the mapper for array type before, such code will always not report any error, and ast print test is correct. Codegen test belongs to the other patch I released. It fits that patch much better.

In D67978#1681388, @ABataev wrote:

In D67978#1681333, @lildmh wrote:

In D67978#1681319, @ABataev wrote:

In D67978#1681318, @lildmh wrote:

Without this patch, it cannot recognize array with mapper, for instance, #pragma omp target map(mapper(a),to: arr[0:2]) won't work without this patch.

What if we have a mapper for the array type?

Without this patch, it won't be able to find the mapper in this case, and Clang assumes no mapper. This patch fixes it, so the compiler can find the mapper in this case.

It means, that you just ignore mappers for the array types, right?

In D67978#1681319, @ABataev wrote:

In D67978#1681318, @lildmh wrote:

Without this patch, it cannot recognize array with mapper, for instance, #pragma omp target map(mapper(a),to: arr[0:2]) won't work without this patch.

What if we have a mapper for the array type?

Without this patch, it cannot recognize array with mapper, for instance, #pragma omp target map(mapper(a),to: arr[0:2]) won't work without this patch.

Fix declare mapper codegen test when the function argument has name attached.

Hi David,

@Meinersbur Hi Michael, could you help me commit this patch? Thanks!

Add a test

In D60972#1607108, @Hahnfeld wrote:

You need to declare them yourself in the test. That's not really elegant, but many other runtime tests already do.

In D60972#1606990, @grokos wrote:

In D60972#1606804, @lildmh wrote:

Since the mapper is not really implemented in this patch, if I add a test, it will be something like below:

__tgt_push_mapper_component(h, base0, begin0, size0, type0);
__tgt_push_mapper_component(h, base1, begin1, size1, type1);
auto total_size = __tgt_mapper_num_components(h);
printf("size=%d", total_size);
// CHECK: size=2

It seems to me this test is not meaningful. I can add a more meaningful test after all mapper patches are upstreamed.
Do you think we need a meaningless test like this now?

Yes, it's good to have a test, even a very elementary one. When full support for declare mapper is upstreamed we can revisit the test and extend it to check real-use scenarios.

Change mapper function argument checking

In D60972#1605250, @grokos wrote:

In D60972#1605210, @lildmh wrote:

In D60972#1605188, @grokos wrote:

Now that the clang-side patch is finalized, this is almost good to go. What is missing from this patch is a test. Can you add something under libomptarget/test/offloading/?

Sure, since this is very simple I'll add a simple test. Do you think the test should be in libomptarget/test/offloading/ or libomptarget/test/api/?

Sorry, I meant libomptarget/test/mapping/!

The api directory is meant for test cases calling omp_* API functions, which is not the case here.

In D60972#1605188, @grokos wrote:

Now that the clang-side patch is finalized, this is almost good to go. What is missing from this patch is a test. Can you add something under libomptarget/test/offloading/?

Thanks Alexey! Could you look into the runtime patch D60972 then?

Ping and rebase

Make emitUDMapperArrayInitOrDel private