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

In D60972#1536070, @Hahnfeld wrote:

In D60972#1533013, @lildmh wrote:

@Hahnfeld Do you really think it is necessary to pass these two functions as arguments, instead of exporting them. If you do, could you explain why?

I don't say it's necessary, but rather an option. Asked differently, why do we need to export them (and decide on externally visible names) if they're only used in a function where we can pass them as arguments?

At the moment, this patch doesn't really add any value (meaning changed behavior) to the runtime. It's needed for D59474 which is already large enough on its own, but it doesn't get to any state where it can be used.

Use tgt instead of kmpc for mapper runtime function names

@Hahnfeld Do you really think it is necessary to pass these two functions as arguments, instead of exporting them. If you do, could you explain why?

Use tgt instead of kmpc

In D60972#1530490, @Hahnfeld wrote:

In D60972#1529265, @lildmh wrote:

In D60972#1529238, @Hahnfeld wrote:

In D60972#1529215, @lildmh wrote:

The compiler doesn't generate code related to std::vector. It's only used in the runtime implementation, so it should be okay with Fortran. Again, the IBM and Intel compiler people seem to agree with it.

Maybe I don't understand where rt_mapper_handle comes from. According to the design slides and D59474, it passed as an argument to the generated omp_mapper[...] function, but how is the runtime system involved in its creation? Will there be additional interface functions / changes that will call this?

The idea is the runtime will create a MapperComponentsTy (std::vector) before calling the mapper function, in, for instance, __tgt_target_data_begin (These parts will be implemented in later patches). When the mapper function is called, the pointer of MapperComponentsTy is passed to it, as void *rt_mapper_handle. The mapper function will call __kmpc_push_mapper_component using this rt_mapper_handle, and then the runtime can put it into the MapperComponentsTy

That's hard to guess from the current patch and isn't in the design slides either (or I have overlooked it). Such information should be made available to reviewers such that they can assess the implementation!

In D60972#1529238, @Hahnfeld wrote:

In D60972#1529215, @lildmh wrote:

In D60972#1529151, @Hahnfeld wrote:

In D60972#1529083, @lildmh wrote:

In D60972#1528662, @Hahnfeld wrote:

It would be great to have such things in public...

Sure, there is no secret. Please see it here if you are interested: https://github.com/lingda-li/public-sharing/blob/master/mapper_runtime_design.pptx

From a quick look, I'd say this does not reflect the current design: The types are named differently, have a different layout (SoA vs AoS) and there's no implementation of __tgt_target_mapper in this patch as @grokos mentioned.

Hi Jonas, starting from slide 8 is the current design, __tgt_target_mapper etc. are deprecated. It may not accurately reflect the current code but the framework should be the same.

Moreover, I'd question the following things:

1. Why are we back to __kmpc_? naming? Most other functions specific to libomptarget are called __tgt_?.

There are other functions in libomptarget starting with __kmpc_, for example, __kmpc_push_target_tripcount.
My understanding is anything that does not directly call the device starting with __kmpc_. The IBM and Intel compiler people seem to be okay with this naming.

Yes, this is the only function AFAICS and it has a comment "will be revised". All other functions related to mapping start with __tgt so unless there is good incentive, we should follow this naming convention.

In D60972#1529151, @Hahnfeld wrote:

In D60972#1529083, @lildmh wrote:

In D60972#1528662, @Hahnfeld wrote:

It would be great to have such things in public...

Sure, there is no secret. Please see it here if you are interested: https://github.com/lingda-li/public-sharing/blob/master/mapper_runtime_design.pptx

From a quick look, I'd say this does not reflect the current design: The types are named differently, have a different layout (SoA vs AoS) and there's no implementation of __tgt_target_mapper in this patch as @grokos mentioned.

Address Alexey's comment about mapping between function and user-defined mappers

In D60972#1528662, @Hahnfeld wrote:

In D60972#1528060, @lildmh wrote:

In D60972#1528023, @grokos wrote:

I think this patch is wrong. Where are the previous additions of target_data and __tgt_target_data_mapper? Possibly you uploaded the diff between the previous revision and the current one, whereas you need to upload the diff between origin/master and your current HEAD!

Please see the attachment of Ravi sent for the biweekly meeting to see details about this new scheme. Sorry I forgot to mention

It would be great to have such things in public...

In D60972#1528023, @grokos wrote:

I think this patch is wrong. Where are the previous additions of target_data and __tgt_target_data_mapper? Possibly you uploaded the diff between the previous revision and the current one, whereas you need to upload the diff between origin/master and your current HEAD!

In D60972#1528023, @grokos wrote:

I think this patch is wrong. Where are the previous additions of target_data and __tgt_target_data_mapper? Possibly you uploaded the diff between the previous revision and the current one, whereas you need to upload the diff between origin/master and your current HEAD!

Implement the new mapper codegen scheme

Implement the new mapper scheme.

Hi Alexey,

Fix code format

Hi Alexey,

Combine 2 pointers into one.

Hi George,