Page MenuHomePhabricator

lildmh (Lingda Li)
User

Projects

User does not belong to any projects.

User Details

User Since
Jan 18 2017, 9:43 AM (191 w, 4 d)

Recent Activity

Aug 19 2020

lildmh added a comment to D84470: [OpenMP 5.0] Fix user-defined mapper privatization in tasks.

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

Aug 19 2020, 8:23 AM · Restricted Project

Jul 27 2020

lildmh added inline comments to D84470: [OpenMP 5.0] Fix user-defined mapper privatization in tasks.
Jul 27 2020, 10:51 AM · Restricted Project

Jul 26 2020

lildmh added a comment to D68100: [OpenMP 5.0] declare mapper runtime implementation.

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];
}
Jul 26 2020, 11:30 AM · Restricted Project

Jul 25 2020

lildmh added a comment to D68100: [OpenMP 5.0] declare mapper runtime implementation.

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?

Jul 25 2020, 8:07 AM · Restricted Project

Jul 24 2020

lildmh added a comment to D68100: [OpenMP 5.0] declare mapper runtime implementation.

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.

Jul 24 2020, 7:26 PM · Restricted Project
lildmh added a comment to D68100: [OpenMP 5.0] declare mapper runtime implementation.

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.

Jul 24 2020, 4:43 PM · Restricted Project
lildmh added a reviewer for D84470: [OpenMP 5.0] Fix user-defined mapper privatization in tasks: ABataev.
Jul 24 2020, 3:26 PM · Restricted Project
lildmh added a comment to D68100: [OpenMP 5.0] declare mapper runtime implementation.

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

Jul 24 2020, 3:24 PM · Restricted Project
lildmh added a comment to D84470: [OpenMP 5.0] Fix user-defined mapper privatization in tasks.

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.

Jul 24 2020, 8:29 AM · Restricted Project

Jul 23 2020

lildmh added a comment to D67833: [OpenMP 5.0] Codegen support to pass user-defined mapper functions to runtime.

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.

Jul 23 2020, 3:52 PM · Restricted Project, Restricted Project, Restricted Project
Herald added a reviewer for D84470: [OpenMP 5.0] Fix user-defined mapper privatization in tasks: jdoerfert.
Jul 23 2020, 3:51 PM · Restricted Project

Jul 22 2020

lildmh added a comment to D68100: [OpenMP 5.0] declare mapper runtime implementation.

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

Jul 22 2020, 6:41 PM · Restricted Project
lildmh added a comment to D67833: [OpenMP 5.0] Codegen support to pass user-defined mapper functions to runtime.

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?

Jul 22 2020, 11:54 AM · Restricted Project, Restricted Project, Restricted Project
lildmh added a comment to D68100: [OpenMP 5.0] declare mapper runtime implementation.

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

Jul 22 2020, 11:47 AM · Restricted Project
lildmh added a comment to D68100: [OpenMP 5.0] declare mapper runtime implementation.

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.

Jul 22 2020, 11:22 AM · Restricted Project
lildmh added a comment to D68100: [OpenMP 5.0] declare mapper runtime implementation.

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

Jul 22 2020, 10:14 AM · Restricted Project
lildmh added a comment to D68100: [OpenMP 5.0] declare mapper runtime implementation.

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

Jul 22 2020, 7:36 AM · Restricted Project
lildmh added a comment to D67833: [OpenMP 5.0] Codegen support to pass user-defined mapper functions to runtime.

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
Jul 22 2020, 6:57 AM · Restricted Project, Restricted Project, Restricted Project

Jul 20 2020

lildmh added a comment to D67833: [OpenMP 5.0] Codegen support to pass user-defined mapper functions to runtime.

Seems that it already has been applied ;-)

Jul 20 2020, 6:53 AM · Restricted Project, Restricted Project, Restricted Project

Jul 16 2020

lildmh added a comment to D83959: Fix compiling warnings in OpenMP declare mapper codegen.

Thanks George!

Jul 16 2020, 11:10 AM · Restricted Project
lildmh added a comment to D67833: [OpenMP 5.0] Codegen support to pass user-defined mapper functions to runtime.

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

Jul 16 2020, 8:50 AM · Restricted Project, Restricted Project, Restricted Project
Herald added a reviewer for D83959: Fix compiling warnings in OpenMP declare mapper codegen: jdoerfert.
Jul 16 2020, 8:48 AM · Restricted Project
lildmh added a comment to D67833: [OpenMP 5.0] Codegen support to pass user-defined mapper functions to runtime.

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.
Jul 16 2020, 6:47 AM · Restricted Project, Restricted Project, Restricted Project

Jul 15 2020

lildmh added a comment to D67833: [OpenMP 5.0] Codegen support to pass user-defined mapper functions to runtime.

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?

Jul 15 2020, 6:02 PM · Restricted Project, Restricted Project, Restricted Project
lildmh updated the diff for D67833: [OpenMP 5.0] Codegen support to pass user-defined mapper functions to runtime.

Include the llvm part

Jul 15 2020, 6:01 PM · Restricted Project, Restricted Project, Restricted Project
lildmh added a comment to D67833: [OpenMP 5.0] Codegen support to pass user-defined mapper functions to runtime.

@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

Jul 15 2020, 1:27 PM · Restricted Project, Restricted Project, Restricted Project
lildmh updated the diff for D67833: [OpenMP 5.0] Codegen support to pass user-defined mapper functions to runtime.

Update diff and pass test. Please accept

Jul 15 2020, 1:13 PM · Restricted Project, Restricted Project, Restricted Project

Jul 14 2020

lildmh requested review of D67833: [OpenMP 5.0] Codegen support to pass user-defined mapper functions to runtime.

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

Jul 14 2020, 1:28 PM · Restricted Project, Restricted Project, Restricted Project
lildmh abandoned D71782: [OpenMP] [NFC] Refactor data mapping to prepare for declare mapper implementation.
Jul 14 2020, 1:04 PM
lildmh accepted D67833: [OpenMP 5.0] Codegen support to pass user-defined mapper functions to runtime.
Jul 14 2020, 12:59 PM · Restricted Project, Restricted Project, Restricted Project
lildmh added a comment to D68100: [OpenMP 5.0] declare mapper runtime implementation.

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

Jul 14 2020, 12:58 PM · Restricted Project
lildmh added a comment to D68100: [OpenMP 5.0] declare mapper runtime implementation.

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

Jul 14 2020, 12:50 PM · Restricted Project
lildmh added a comment to D68100: [OpenMP 5.0] declare mapper runtime implementation.

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

Jul 14 2020, 12:48 PM · Restricted Project

Mar 30 2020

lildmh added inline comments to D77005: [OpenMP] Optimized stream selection by scheduling data mapping for the same target region into a same stream.
Mar 30 2020, 5:55 AM · Restricted Project

Feb 14 2020

lildmh added a comment to D68100: [OpenMP 5.0] declare mapper runtime implementation.

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.

Feb 14 2020, 11:07 AM · Restricted Project

Dec 23 2019

lildmh updated the diff for D68100: [OpenMP 5.0] declare mapper runtime implementation.

Address Jon's comments

Dec 23 2019, 7:44 AM · Restricted Project
lildmh added inline comments to D68100: [OpenMP 5.0] declare mapper runtime implementation.
Dec 23 2019, 7:44 AM · Restricted Project
lildmh updated the diff for D68100: [OpenMP 5.0] declare mapper runtime implementation.

Rebase and rediff with the nfc version

Dec 23 2019, 6:49 AM · Restricted Project

Dec 21 2019

lildmh added a comment to D71782: [OpenMP] [NFC] Refactor data mapping to prepare for declare mapper implementation.

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

Dec 21 2019, 2:36 PM

Dec 20 2019

lildmh created D71782: [OpenMP] [NFC] Refactor data mapping to prepare for declare mapper implementation.
Dec 20 2019, 1:25 PM

Dec 18 2019

lildmh added inline comments to D67833: [OpenMP 5.0] Codegen support to pass user-defined mapper functions to runtime.
Dec 18 2019, 11:35 AM · Restricted Project, Restricted Project, Restricted Project
lildmh added inline comments to D67833: [OpenMP 5.0] Codegen support to pass user-defined mapper functions to runtime.
Dec 18 2019, 11:16 AM · Restricted Project, Restricted Project, Restricted Project
lildmh added inline comments to D67833: [OpenMP 5.0] Codegen support to pass user-defined mapper functions to runtime.
Dec 18 2019, 10:56 AM · Restricted Project, Restricted Project, Restricted Project
lildmh updated the diff for D67833: [OpenMP 5.0] Codegen support to pass user-defined mapper functions to runtime.

Change the function name

Dec 18 2019, 10:45 AM · Restricted Project, Restricted Project, Restricted Project
lildmh added inline comments to D67833: [OpenMP 5.0] Codegen support to pass user-defined mapper functions to runtime.
Dec 18 2019, 10:35 AM · Restricted Project, Restricted Project, Restricted Project
lildmh added inline comments to D67833: [OpenMP 5.0] Codegen support to pass user-defined mapper functions to runtime.
Dec 18 2019, 9:56 AM · Restricted Project, Restricted Project, Restricted Project
lildmh added inline comments to D67833: [OpenMP 5.0] Codegen support to pass user-defined mapper functions to runtime.
Dec 18 2019, 9:41 AM · Restricted Project, Restricted Project, Restricted Project
lildmh updated the diff for D67833: [OpenMP 5.0] Codegen support to pass user-defined mapper functions to runtime.

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

Dec 18 2019, 8:53 AM · Restricted Project, Restricted Project, Restricted Project

Dec 17 2019

lildmh added inline comments to D67833: [OpenMP 5.0] Codegen support to pass user-defined mapper functions to runtime.
Dec 17 2019, 7:49 AM · Restricted Project, Restricted Project, Restricted Project

Dec 16 2019

lildmh added inline comments to D67833: [OpenMP 5.0] Codegen support to pass user-defined mapper functions to runtime.
Dec 16 2019, 4:44 PM · Restricted Project, Restricted Project, Restricted Project
lildmh added a comment to D67833: [OpenMP 5.0] Codegen support to pass user-defined mapper functions to runtime.

Alexey, thanks for the review

Dec 16 2019, 12:43 PM · Restricted Project, Restricted Project, Restricted Project
lildmh added a comment to D67833: [OpenMP 5.0] Codegen support to pass user-defined mapper functions to runtime.

I'll split the patch into 2 later

Dec 16 2019, 11:40 AM · Restricted Project, Restricted Project, Restricted Project
lildmh updated the diff for D67833: [OpenMP 5.0] Codegen support to pass user-defined mapper functions to runtime.

Rebase and address comments

Dec 16 2019, 11:40 AM · Restricted Project, Restricted Project, Restricted Project

Nov 25 2019

lildmh updated the diff for D68100: [OpenMP 5.0] declare mapper runtime implementation.

Thanks for your reviews. Hope this looks better.

Nov 25 2019, 2:46 PM · Restricted Project

Nov 15 2019

lildmh added inline comments to D68100: [OpenMP 5.0] declare mapper runtime implementation.
Nov 15 2019, 8:54 AM · Restricted Project

Nov 13 2019

lildmh updated the diff for D67833: [OpenMP 5.0] Codegen support to pass user-defined mapper functions to runtime.

Rebase

Nov 13 2019, 4:22 PM · Restricted Project, Restricted Project, Restricted Project
lildmh added inline comments to D68100: [OpenMP 5.0] declare mapper runtime implementation.
Nov 13 2019, 4:13 PM · Restricted Project
lildmh updated the diff for D68100: [OpenMP 5.0] declare mapper runtime implementation.

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

Nov 13 2019, 4:13 PM · Restricted Project

Sep 27 2019

lildmh added a comment to D68100: [OpenMP 5.0] declare mapper runtime implementation.

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.

Sep 27 2019, 1:33 PM · Restricted Project
lildmh added a comment to D68100: [OpenMP 5.0] declare mapper runtime implementation.

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.

Sep 27 2019, 1:33 PM · Restricted Project
lildmh added a comment to D68100: [OpenMP 5.0] declare mapper runtime implementation.

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.

Sep 27 2019, 1:14 PM · Restricted Project
lildmh added a comment to D68100: [OpenMP 5.0] declare mapper runtime implementation.

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

Sep 27 2019, 1:02 PM · Restricted Project
lildmh added a comment to D68100: [OpenMP 5.0] declare mapper runtime implementation.

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

Sep 27 2019, 12:46 PM · Restricted Project
lildmh added a comment to D68100: [OpenMP 5.0] declare mapper runtime implementation.

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

Sep 27 2019, 12:43 PM · Restricted Project
lildmh added a comment to D68100: [OpenMP 5.0] declare mapper runtime implementation.

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.

Sep 27 2019, 12:17 PM · Restricted Project
lildmh updated the diff for D68100: [OpenMP 5.0] declare mapper runtime implementation.

Please review when you have time, thanks

Sep 27 2019, 7:44 AM · Restricted Project

Sep 26 2019

lildmh updated the diff for D67833: [OpenMP 5.0] Codegen support to pass user-defined mapper functions to runtime.

Rebase

Sep 26 2019, 6:08 PM · Restricted Project, Restricted Project, Restricted Project
lildmh added a comment to D67978: [OpenMP 5.0] Fix user-defined mapper lookup in sema.

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.

Sep 26 2019, 1:42 PM · Restricted Project, Restricted Project, Restricted Project
lildmh added inline comments to D67978: [OpenMP 5.0] Fix user-defined mapper lookup in sema.
Sep 26 2019, 1:31 PM · Restricted Project, Restricted Project, Restricted Project
lildmh added inline comments to D67978: [OpenMP 5.0] Fix user-defined mapper lookup in sema.
Sep 26 2019, 1:13 PM · Restricted Project, Restricted Project, Restricted Project
lildmh created D68100: [OpenMP 5.0] declare mapper runtime implementation.
Sep 26 2019, 12:34 PM · Restricted Project
lildmh added inline comments to D67978: [OpenMP 5.0] Fix user-defined mapper lookup in sema.
Sep 26 2019, 11:55 AM · Restricted Project, Restricted Project, Restricted Project
lildmh added inline comments to D67978: [OpenMP 5.0] Fix user-defined mapper lookup in sema.
Sep 26 2019, 11:32 AM · Restricted Project, Restricted Project, Restricted Project
lildmh updated the diff for D67978: [OpenMP 5.0] Fix user-defined mapper lookup in sema.

Fix mapper type checking

Sep 26 2019, 11:17 AM · Restricted Project, Restricted Project, Restricted Project
lildmh updated the diff for D67978: [OpenMP 5.0] Fix user-defined mapper lookup in sema.
Sep 26 2019, 11:17 AM · Restricted Project, Restricted Project, Restricted Project
lildmh added a comment to D67978: [OpenMP 5.0] Fix user-defined mapper lookup in sema.

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.

Sep 26 2019, 9:21 AM · Restricted Project, Restricted Project, Restricted Project
lildmh added a comment to D67978: [OpenMP 5.0] Fix user-defined mapper lookup in sema.

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.

Sep 26 2019, 6:11 AM · Restricted Project, Restricted Project, Restricted Project

Sep 25 2019

lildmh added a comment to D67978: [OpenMP 5.0] Fix user-defined mapper lookup in sema.

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.

Sep 25 2019, 1:55 PM · Restricted Project, Restricted Project, Restricted Project

Sep 24 2019

lildmh added a comment to D67978: [OpenMP 5.0] Fix user-defined mapper lookup in sema.

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?

Sep 24 2019, 5:04 PM · Restricted Project, Restricted Project, Restricted Project
lildmh added a comment to D67978: [OpenMP 5.0] Fix user-defined mapper lookup in sema.

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?

Sep 24 2019, 1:58 PM · Restricted Project, Restricted Project, Restricted Project
lildmh added a comment to D67978: [OpenMP 5.0] Fix user-defined mapper lookup in sema.

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.

Sep 24 2019, 1:50 PM · Restricted Project, Restricted Project, Restricted Project
lildmh updated the diff for D67833: [OpenMP 5.0] Codegen support to pass user-defined mapper functions to runtime.
Sep 24 2019, 1:04 PM · Restricted Project, Restricted Project, Restricted Project
lildmh added reviewers for D67833: [OpenMP 5.0] Codegen support to pass user-defined mapper functions to runtime: ABataev, Meinersbur, hfinkel.
Sep 24 2019, 1:04 PM · Restricted Project, Restricted Project, Restricted Project
lildmh created D67978: [OpenMP 5.0] Fix user-defined mapper lookup in sema.
Sep 24 2019, 12:59 PM · Restricted Project, Restricted Project, Restricted Project

Sep 20 2019

lildmh created D67833: [OpenMP 5.0] Codegen support to pass user-defined mapper functions to runtime.
Sep 20 2019, 6:54 AM · Restricted Project, Restricted Project, Restricted Project

Aug 5 2019

lildmh updated the diff for D59474: [OpenMP 5.0] Codegen support for user-defined mappers.

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

Aug 5 2019, 5:57 AM · Restricted Project, Restricted Project, Restricted Project

Aug 2 2019

lildmh added a comment to D56326: [OpenMP 5.0] Parsing/sema support for "omp declare mapper" directive.

Hi David,

Aug 2 2019, 1:10 PM · Restricted Project, Restricted Project

Aug 1 2019

lildmh added a comment to D60972: [OpenMP 5.0] libomptarget interface for declare mapper functions.

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

Aug 1 2019, 6:06 AM · Restricted Project, Restricted Project

Jul 31 2019

lildmh updated the diff for D60972: [OpenMP 5.0] libomptarget interface for declare mapper functions.

Add a test

Jul 31 2019, 7:39 AM · Restricted Project, Restricted Project

Jul 30 2019

lildmh added a comment to D60972: [OpenMP 5.0] libomptarget interface for declare mapper functions.

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

Jul 30 2019, 1:46 PM · Restricted Project, Restricted Project
lildmh added a comment to D60972: [OpenMP 5.0] libomptarget interface for declare mapper functions.

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.

Jul 30 2019, 1:15 PM · Restricted Project, Restricted Project
lildmh updated the diff for D59474: [OpenMP 5.0] Codegen support for user-defined mappers.

Change mapper function argument checking

Jul 30 2019, 11:45 AM · Restricted Project, Restricted Project, Restricted Project
lildmh added a comment to D60972: [OpenMP 5.0] libomptarget interface for declare mapper functions.

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.

Jul 30 2019, 11:36 AM · Restricted Project, Restricted Project

Jul 29 2019

lildmh added a comment to D60972: [OpenMP 5.0] libomptarget interface for declare mapper functions.

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

Jul 29 2019, 2:06 PM · Restricted Project, Restricted Project
lildmh added inline comments to D59474: [OpenMP 5.0] Codegen support for user-defined mappers.
Jul 29 2019, 1:45 PM · Restricted Project, Restricted Project, Restricted Project
lildmh added inline comments to D59474: [OpenMP 5.0] Codegen support for user-defined mappers.
Jul 29 2019, 1:19 PM · Restricted Project, Restricted Project, Restricted Project
lildmh added a comment to D59474: [OpenMP 5.0] Codegen support for user-defined mappers.

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

Jul 29 2019, 12:10 PM · Restricted Project, Restricted Project, Restricted Project
lildmh updated the diff for D60972: [OpenMP 5.0] libomptarget interface for declare mapper functions.

Ping and rebase

Jul 29 2019, 12:08 PM · Restricted Project, Restricted Project

Jul 26 2019

lildmh updated the diff for D59474: [OpenMP 5.0] Codegen support for user-defined mappers.

Make emitUDMapperArrayInitOrDel private

Jul 26 2019, 12:52 PM · Restricted Project, Restricted Project, Restricted Project
lildmh added inline comments to D59474: [OpenMP 5.0] Codegen support for user-defined mappers.
Jul 26 2019, 11:04 AM · Restricted Project, Restricted Project, Restricted Project