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

Remove virtual from function declaration

Get rid of MSVC requirement of this, and a virtual function

Rebase

In D59474#1589464, @ABataev wrote:

In D59474#1561161, @lildmh wrote:

Change the type of size from size_t to int64_t, and rebase

Lingda, could you rebase it again? Thanks.

Change the type of size from size_t to int64_t, and rebase

Address Alexey's comments

Fix mapper function name mangling