In D78075#2272474, @tianshilei1992 wrote:

In D78075#2272398, @ye-luo wrote:

However, OpenMP task has a problem that it must be within
to a parallel region; otherwise the task will be executed immediately. As a
result, if we directly wrap to a regular task, the nowait target outside of a
parallel region is still a synchronous version.

The spec says an implicit task can be generated by an implicit parallel region which can be the whole OpenMP program. For this reason, the need of explicit parallel region is a limitation of the llvm OpenMP runtime, right?

Can I have an option to run the nowait region as a regular task instead of an unshackled task? So I can use "parallel" and well established ways to control the thread affinity.

According to the spec, an implicit parallel region is an inactive parallel region that is not generated from a parallel construct. And based on the definition of active parallel region, which is a parallel region that is executed by a team consisting of more than one thread, an inactive parallel region only has one thread. Since we only have one thread, if we encounter a task, executing it immediately does make sense as we don't have another thread to execute it.

If I remember correctly, you may yield the thread inside a target region after enqueuing kernels and transfers. So even with 1 thread, there is chance to run other tasks without finishing this target. Isn't that possible?

However, OpenMP task has a problem that it must be within
to a parallel region; otherwise the task will be executed immediately. As a
result, if we directly wrap to a regular task, the nowait target outside of a
parallel region is still a synchronous version.

The changes I requested as been added. Remove my blocking. Still need other reviews to be addressed.

In D87165#2257164, @jhuber6 wrote:

Added additional comments. Should I add them to the doxygen notes at the top?

In D86804#2248214, @jdoerfert wrote:

In D86804#2248201, @ye-luo wrote:

In D86804#2247785, @jdoerfert wrote:

In D86804#2247697, @ye-luo wrote:

It seems that functions are marked static so they should be OK. However, including the whole Debug.h in a plugin cpp makes it feel OK to use any function/macro from the header file. But actually only part of the macros are for the plugin. some are only for the libomptarget.

I'm not sure we want to make a distinction, the point is to move to a unified debug/message model. You can choose not only the level of information but also the kind of output (text, json, ...). The messages will then be tied to the webpage via enums, that allows all plugins to emit the same message for the same thing with the same link to more information. There will certainly things that are only used in libomptarget or the plugins, but I don't see how that is any worse than duplicating the parts that are used by both.

I didn't mean to duplicate anything. Instead you need multiple header files. One for common parts, one for libomptarget and one for plugins. The latter two both include the first one. Later you expand OFFLOAD_XXX signals, they can be added to the common file. The return signal is generated by the plugins and captured by the libomptarget. Some users may want to see only the messages captured universally by libomptarget. Some users still want to see the native error message. So the libomptarget and plugin side error handling still needs to be separated.

I fail to see why this machinery is necessary to emit only messages from one place and not the other. I am not against a hierarchy of headers per se, but right now, and maybe also later, there seems to be little point.
I mean, we need to introduce a new env variable, actually two, that allow separate control. Once we have that we can argue about separation.
Alternatively, I would have suggested to define the "location" prior to including the debug header, e.g.,:

#define DEBUG_LOCATION "PLUGIN"
#include "Debug.h"

which we verify like:

#if DEBUG_LOCATION != "PLUGIN" and DEBUG_LOCATION != "OMPTARGET"
#error ...
#endif

At the end of the day I want to simplify things. A single location for all our debug needs sounds simpler than 1 + #plugins to me, even if we don't use all functionality at each location. If separation does not allow anything we cannot reasonably do in a single location, I doubt it provides a benefit.

In D86804#2247785, @jdoerfert wrote:

In D86804#2247697, @ye-luo wrote:

It seems that functions are marked static so they should be OK. However, including the whole Debug.h in a plugin cpp makes it feel OK to use any function/macro from the header file. But actually only part of the macros are for the plugin. some are only for the libomptarget.

I'm not sure we want to make a distinction, the point is to move to a unified debug/message model. You can choose not only the level of information but also the kind of output (text, json, ...). The messages will then be tied to the webpage via enums, that allows all plugins to emit the same message for the same thing with the same link to more information. There will certainly things that are only used in libomptarget or the plugins, but I don't see how that is any worse than duplicating the parts that are used by both.

It seems that functions are marked static so they should be OK. However, including the whole Debug.h in a plugin cpp makes it feel OK to use any function/macro from the header file. But actually only part of the macros are for the plugin. some are only for the libomptarget.

I don't feel right having Debug.h shared by libomptarget and plugins especially when Debug.h is not just macro but also functions.

Please document the flags in the patch summary.

I prefer to PrivateArgumentManagerTy moved into its own files.
The rest looks good to me.

Down the road, we may need a way to allocate host pinned memory via the plugin for the host buffer to maximize transfer performance.

In D81054#2229637, @JonChesterfield wrote:

As a heads up, I'm told this breaks amdgpu tests. @ronlieb is looking at the merge from upstream, don't have any more details at this time. The basic idea of wrapping device alloc seems likely to be sound for all targets so I'd guess we've run into a bug in this patch.

Only minor things.

In D86307#2229059, @jdoerfert wrote:

In D86307#2229032, @ye-luo wrote:

Why just "small" ones? why not all of them?

In addition to the last paragraph of the new commit message, we also have to copy the data on the host in the right place. That is not free as the size grows.

Why just "small" ones? why not all of them?

LGTM

LGTM

In addition,

1. the DeviceTy copy constructor and assign operator are imperfect before this patch. I don't think we can fix them in this patch. We should just document the imperfection here.
2. Because the memory limit is per allocation, it seems that the MemoryManager can still hold infinite amount of memory and we don't have way to free them. I'm concerned about having this feature on by default.

What is the current status of this patch?
@lildmh could you update this patch? I'd like to test it against
https://bugs.llvm.org/show_bug.cgi?id=47122

Block the patch temporarily for my earlier questions.

1. Please mention LIBOMPTARGET_MEMORY_MANAGER_THRESHOLD, default value and unit in the patch summary.
2. Is it possible to have a unit test testing the manager class behaviors?
3. Can we offload to host and run address sanitizer or valgrind?

I'm not sure if I'm asking for too much here.

LGTM

Thanks for fixing the bug. It should be good for the moment.
When I think about the existence of recursive mapper, we may still have more sync than needed. I think recursion the whole targetDataBegin/targetDataEnd is convenient but sub-optimal choice.
Recursion should only be done on the map/mapper analysis. Just leave my thoughts here. It needs a discussion beyond this patch.

LGTM.

LGTM. Please mention renaming variables in the summary.

LGTM. My applications run as expected now. PR46824, PR46012, PR46868 all work fine.

Only minor documentation issues.

In D84767#2180280, @ye-luo wrote:

This patch
GPU activities: 96.99% 350.05ms 10 35.005ms 1.5680us 350.00ms [CUDA memcpy HtoD]
before the July21 change
GPU activities: 95.33% 20.317ms 4 5.0793ms 1.6000us 20.305ms [CUDA memcpy HtoD]
Still more transfer than it should.

This patch
GPU activities: 96.99% 350.05ms 10 35.005ms 1.5680us 350.00ms [CUDA memcpy HtoD]
before the July21 change
GPU activities: 95.33% 20.317ms 4 5.0793ms 1.6000us 20.305ms [CUDA memcpy HtoD]
Still more transfer than it should.

LGTM

OK. Leave the unrelated renaming to the future.

Only one minor issue. Your initial sophisticated patch made my thought you replaced all the lock/unlock. After splitting, the change becomes very clean.

Should be easy to address my comments and let us get this merged ASAP.

In D84778#2179627, @tianshilei1992 wrote:

I don't think it deserves three patches. The goal is to refactor the target function, and this patch just did this only thing. According to the bi-weekly meeting, the renaming could be with other related changes.

In D84778#2179502, @ye-luo wrote:

In addtion, should we update target_data_update as well?

I didn't touch target_data_update. Basically I only take care of related code.

Please check the reproducer in https://bugs.llvm.org/show_bug.cgi?id=46868 with LIBOMPTARGET_DEBUG=1.
The reference counting on the base pointer variable has side effects. It was not cleaned up when these variables leave its scope.

Needs to split this patch into three.

1. function renaming. In addtion, should we update target_data_update as well?
2. std::lock_guard change.
3. "target" change.

The order of 1 and 2 can be flexible

In D84182#2173578, @grokos wrote:

@ABataev:

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

#include <stdio.h>

int *yptr;

int main() {
  int y[10];
  y[1] = 1;
  yptr = &y[0];

  printf("&yptr = %p\n", &yptr);
  printf("&y[0] = %p\n", &y[0]);

  #pragma omp target data map(to: yptr[0:5])
  #pragma omp target
  {
    printf("y = %d\n", yptr[1]);
    yptr[1] = 10;
    printf("y = %d\n", yptr[1]);
  }

  printf("y = %d\n", yptr[1]);
  return 0;
}

The arguments clang generates are:

1) base = &y[0], begin = &yptr, size = 8, type = TARGET_PARAM | TO
2) base = &yptr, begin = &y[0], size = 8, type = PTR_AND_OBJ | TO

The second argument is correct, the first argument doesn't make much sense. I believe it should have its base set to &yptr, not &y[0].
y[0] is not the base for anything, it's only the pointee object.

In D82074#2173271, @protze.joachim wrote:

The cmake version was bumped to 3.14

OK. It is less broken now.
target_data_end still does Device.deallocTgtPtr and needs a sync before it.
To fully fix this issue, target_data_end must be spitted.

Indeed, target_data_begin should be split as well. cudaMalloc blocks the whole device. Alternating cudaMalloc and transfer only makes the whole process further slower. Better to make all the allocation and then start queuing the transfer.

Does it mean the D2H will always run synchronously after this change?
Does it also mean that target_data_end should be split into data transfer and data free parts?

I verified that 46012 is fixed with this patch

Could you describe what "SplitMemTransfer" is? The current patch summary doesn't provide sufficient explanation. What kind of offloading arrays is considered by this optimization?

In D83963#2156786, @jdoerfert wrote:

Add flush

In D83963#2156495, @grokos wrote:

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

Are you sure abort gets fprintf flushed?

LGTM

Assertion failure in PR46687 goes away with this patch.

@protze.joachim Is the current way of handling CU_CTX_SCHED_BLOCKING_SYNC OK?
I mean still set the flag as CU_CTX_SCHED_BLOCKING_SYNC if the primary context is not activated prior to libomptarget.
Allowing users to choose a scheme will be a separate patch.
I'd like to get this patch in as soon as possible because the 11 release branch will be created soon and the current issue is a stopper for applications using libomptarget.

Updated description regarding setting CU_CTX_SCHED_BLOCKING_SYNC.

In D82718#2123747, @tianshilei1992 wrote:

In D82718#2123675, @protze.joachim wrote:

This patch drops the CU_CTX_SCHED_BLOCKING_SYNC property currently selected for the context. Is this intended? Should we add another function call to request this behavior for the primary context?

That is good point. We depend on the synchronous behavior in some cases in the RTL.

In D82718#2123675, @protze.joachim wrote:

This patch drops the CU_CTX_SCHED_BLOCKING_SYNC property currently selected for the context. Is this intended? Should we add another function call to request this behavior for the primary context?

I have updated the patch removing the old behavior. @Hahnfeld would you approve?

It seems no one objects to removing the old behavior. I will update the patch.

In D82789#2120848, @tianshilei1992 wrote:

Thank you for the patch. Just one tiny comment: are these everything for the warning about the debug output?

In D82718#2120164, @jdoerfert wrote:

In D82718#2119143, @ye-luo wrote:

I'm not aware of any side effect but others may have different opinions. Since it is low effort to keep both, I made the choice to keep the old behavior accessible.

More options mean more configurations we should test. Not that we do it right now but we should. Let's see if anyone else has an opinion about this.

In D82718#2120235, @Hahnfeld wrote:

(In general, all patches must be sent to the respective -commits list. This also makes feedback more likely.)