Page MenuHomePhabricator

ye-luo (Ye Luo)
User

Projects

User does not belong to any projects.

User Details

User Since
Aug 8 2018, 8:02 AM (110 w, 3 d)

Recent Activity

Yesterday

ye-luo updated the diff for D87980: [OpenMP] Protect unrecogonized CUDA error code.
Sat, Sep 19, 8:44 PM
ye-luo requested review of D87980: [OpenMP] Protect unrecogonized CUDA error code.
Sat, Sep 19, 8:31 PM

Mon, Sep 14

ye-luo added a comment to D78075: [WIP][Clang][OpenMP] Added support for nowait target in CodeGen.

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?

Mon, Sep 14, 6:11 PM · Restricted Project
ye-luo added a comment to D78075: [WIP][Clang][OpenMP] Added support for nowait target in CodeGen.

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.

Mon, Sep 14, 3:51 PM · Restricted Project

Thu, Sep 10

ye-luo added inline comments to D87165: [OpenMP] Begin Printing Information Dumps In Libomptarget and Plugins.
Thu, Sep 10, 6:23 AM · Restricted Project

Tue, Sep 8

ye-luo accepted D87165: [OpenMP] Begin Printing Information Dumps In Libomptarget and Plugins.

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

Tue, Sep 8, 8:27 AM · Restricted Project

Fri, Sep 4

ye-luo added a comment to D87165: [OpenMP] Begin Printing Information Dumps In Libomptarget and Plugins.

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

Fri, Sep 4, 2:58 PM · Restricted Project
ye-luo requested changes to D87165: [OpenMP] Begin Printing Information Dumps In Libomptarget and Plugins.
Fri, Sep 4, 2:32 PM · Restricted Project

Tue, Sep 1

ye-luo added inline comments to D77609: [OpenMP][WIP] Added the support for unshackled task in RTL.
Tue, Sep 1, 2:00 PM · Restricted Project

Mon, Aug 31

ye-luo added a comment to D86804: [OpenMP] Consolidate error handling and debug messages in Libomptarget.

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.

Mon, Aug 31, 3:22 PM · Restricted Project
ye-luo added a comment to D86804: [OpenMP] Consolidate error handling and debug messages in Libomptarget.
Mon, Aug 31, 3:05 PM · Restricted Project
ye-luo added a comment to D86804: [OpenMP] Consolidate error handling and debug messages in Libomptarget.

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.

Mon, Aug 31, 2:30 PM · Restricted Project
ye-luo added a comment to D86804: [OpenMP] Consolidate error handling and debug messages in Libomptarget.

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.

Mon, Aug 31, 10:10 AM · Restricted Project
ye-luo added a comment to D86804: [OpenMP] Consolidate error handling and debug messages in 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.

Mon, Aug 31, 9:41 AM · Restricted Project

Wed, Aug 26

ye-luo added a comment to D86483: [OpenMP] Always emit debug messages that indicate offloading failure.

Please document the flags in the patch summary.

Wed, Aug 26, 9:48 AM · Restricted Project

Mon, Aug 24

ye-luo accepted D86307: [OpenMP] Pack first-private arguments to improve efficiency of data transfer.

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

Mon, Aug 24, 2:26 PM · Restricted Project
ye-luo added inline comments to D86307: [OpenMP] Pack first-private arguments to improve efficiency of data transfer.
Mon, Aug 24, 10:44 AM · Restricted Project

Sun, Aug 23

ye-luo added a comment to D86307: [OpenMP] Pack first-private arguments to improve efficiency of data transfer.

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

Sun, Aug 23, 9:41 PM · Restricted Project

Aug 20 2020

ye-luo added a comment to D81054: [OpenMP] Introduce target memory manager.

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.

Aug 20 2020, 6:33 PM · Restricted Project
ye-luo added a comment to D86307: [OpenMP] Pack first-private arguments to improve efficiency of data transfer.

Only minor things.

Aug 20 2020, 3:58 PM · Restricted Project
ye-luo added inline comments to D86307: [OpenMP] Pack first-private arguments to improve efficiency of data transfer.
Aug 20 2020, 3:18 PM · Restricted Project
ye-luo added inline comments to D86307: [OpenMP] Pack first-private arguments to improve efficiency of data transfer.
Aug 20 2020, 1:51 PM · Restricted Project
ye-luo added inline comments to D86307: [OpenMP] Pack first-private arguments to improve efficiency of data transfer.
Aug 20 2020, 12:47 PM · Restricted Project
ye-luo added a comment to D86307: [OpenMP] Pack first-private arguments to improve efficiency of data transfer.

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.

Aug 20 2020, 12:12 PM · Restricted Project
ye-luo added a comment to D86307: [OpenMP] Pack first-private arguments to improve efficiency of data transfer.

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

Aug 20 2020, 11:35 AM · Restricted Project

Aug 19 2020

ye-luo accepted D81054: [OpenMP] Introduce target memory manager.

LGTM

Aug 19 2020, 1:18 PM · Restricted Project
ye-luo added inline comments to D81054: [OpenMP] Introduce target memory manager.
Aug 19 2020, 12:34 PM · Restricted Project
ye-luo accepted D86238: [OpenMP] Refactored the function `DeviceTy::data_exchange`.

LGTM

Aug 19 2020, 12:03 PM · Restricted Project

Aug 18 2020

ye-luo added inline comments to D81054: [OpenMP] Introduce target memory manager.
Aug 18 2020, 8:47 PM · Restricted Project
ye-luo added inline comments to D81054: [OpenMP] Introduce target memory manager.
Aug 18 2020, 8:27 PM · Restricted Project
ye-luo added inline comments to D81054: [OpenMP] Introduce target memory manager.
Aug 18 2020, 8:01 PM · Restricted Project
ye-luo added a comment to D81054: [OpenMP] Introduce target memory manager.

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.
Aug 18 2020, 6:57 PM · Restricted Project

Aug 13 2020

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

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

Aug 13 2020, 3:00 PM · Restricted Project

Aug 12 2020

ye-luo added inline comments to D81054: [OpenMP] Introduce target memory manager.
Aug 12 2020, 6:23 PM · Restricted Project
ye-luo added inline comments to D81054: [OpenMP] Introduce target memory manager.
Aug 12 2020, 5:31 PM · Restricted Project
ye-luo added inline comments to D81054: [OpenMP] Introduce target memory manager.
Aug 12 2020, 5:09 PM · Restricted Project
ye-luo added inline comments to D81054: [OpenMP] Introduce target memory manager.
Aug 12 2020, 4:54 PM · Restricted Project
ye-luo added inline comments to D81054: [OpenMP] Introduce target memory manager.
Aug 12 2020, 4:49 PM · Restricted Project
ye-luo added inline comments to D81054: [OpenMP] Introduce target memory manager.
Aug 12 2020, 10:22 AM · Restricted Project
ye-luo requested changes to D81054: [OpenMP] Introduce target memory manager.

Block the patch temporarily for my earlier questions.

Aug 12 2020, 10:11 AM · Restricted Project
ye-luo added a comment to D81054: [OpenMP] Introduce target memory manager.
  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.

Aug 12 2020, 10:05 AM · Restricted Project

Jul 31 2020

ye-luo accepted D84996: [OpenMP] Fixed the issue that target memory deallocation might be called when they're being used.

LGTM

Jul 31 2020, 2:06 PM · Restricted Project

Jul 30 2020

ye-luo added a comment to D84996: [OpenMP] Fixed the issue that target memory deallocation might be called when they're being used.

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.

Jul 30 2020, 7:02 PM · Restricted Project
ye-luo accepted D84816: [OpenMP] Refactored the function `target`.

LGTM.

Jul 30 2020, 5:06 PM · Restricted Project
ye-luo accepted D84991: [OpenMP] Refactored the function `targetDataEnd`.

LGTM. Please mention renaming variables in the summary.

Jul 30 2020, 5:04 PM · Restricted Project

Jul 29 2020

ye-luo accepted D84767: [OPENMP]Fix PR46824: Global declare target pointer cannot be accessed in target region..

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

Jul 29 2020, 2:09 PM · Restricted Project, Restricted Project
ye-luo added a comment to D84816: [OpenMP] Refactored the function `target`.

Only minor documentation issues.

Jul 29 2020, 9:32 AM · Restricted Project

Jul 28 2020

ye-luo added a comment to D84767: [OPENMP]Fix PR46824: Global declare target pointer cannot be accessed in target region..

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.

Jul 28 2020, 4:13 PM · Restricted Project, Restricted Project
ye-luo added a comment to D84767: [OPENMP]Fix PR46824: Global declare target pointer cannot be accessed in target region..

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.

Jul 28 2020, 3:12 PM · Restricted Project, Restricted Project
ye-luo accepted D84799: [OpenMP] Replaced mutex lock/unlock in `target` with `std::lock_guard`.

LGTM

Jul 28 2020, 2:16 PM · Restricted Project
ye-luo accepted D84797: [NFC][OpenMP] Renamed all variable and function names in `target` to conform with LLVM code standard.

OK. Leave the unrelated renaming to the future.

Jul 28 2020, 2:15 PM · Restricted Project
ye-luo added a comment to D84799: [OpenMP] Replaced mutex lock/unlock in `target` with `std::lock_guard`.

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

Jul 28 2020, 2:04 PM · Restricted Project
ye-luo added a comment to D84797: [NFC][OpenMP] Renamed all variable and function names in `target` to conform with LLVM code standard.

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

Jul 28 2020, 1:56 PM · Restricted Project
ye-luo added a comment to D84778: [OpenMP] Refactor the `target` function.

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 addtion, should we update target_data_update as well?

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

Jul 28 2020, 12:57 PM · Restricted Project
ye-luo requested changes to D84767: [OPENMP]Fix PR46824: Global declare target pointer cannot be accessed in target region..

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.

Jul 28 2020, 12:42 PM · Restricted Project, Restricted Project
ye-luo requested changes to D84778: [OpenMP] Refactor the `target` function.

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

Jul 28 2020, 11:30 AM · Restricted Project
ye-luo added a comment to D84182: [OPENMP]Fix PR46012: declare target pointer cannot be accessed in target region..

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

Jul 28 2020, 6:33 AM · Restricted Project

Jul 24 2020

ye-luo added inline comments to D84487: [OpenMP] Add more pass-through functions in DeviceTy.
Jul 24 2020, 3:37 PM · Restricted Project
ye-luo updated the diff for D84487: [OpenMP] Add more pass-through functions in DeviceTy.
Jul 24 2020, 3:36 PM · Restricted Project
ye-luo added a comment to D82074: [OpenMP] Set cmake policies CMP0074 and CMP0075 to NEW.

The cmake version was bumped to 3.14

Jul 24 2020, 2:09 PM · Restricted Project
ye-luo added inline comments to D84487: [OpenMP] Add more pass-through functions in DeviceTy.
Jul 24 2020, 2:06 PM · Restricted Project
ye-luo updated the diff for D84487: [OpenMP] Add more pass-through functions in DeviceTy.
Jul 24 2020, 2:03 PM · Restricted Project
ye-luo added inline comments to D84487: [OpenMP] Add more pass-through functions in DeviceTy.
Jul 24 2020, 12:34 PM · Restricted Project
ye-luo added inline comments to D84487: [OpenMP] Add more pass-through functions in DeviceTy.
Jul 24 2020, 10:32 AM · Restricted Project
ye-luo added inline comments to D84487: [OpenMP] Add more pass-through functions in DeviceTy.
Jul 24 2020, 9:45 AM · Restricted Project
ye-luo added inline comments to D84487: [OpenMP] Add more pass-through functions in DeviceTy.
Jul 24 2020, 8:51 AM · Restricted Project
ye-luo updated the diff for D84487: [OpenMP] Add more pass-through functions in DeviceTy.
Jul 24 2020, 8:19 AM · Restricted Project
ye-luo added inline comments to D84487: [OpenMP] Add more pass-through functions in DeviceTy.
Jul 24 2020, 7:28 AM · Restricted Project
ye-luo added inline comments to D84487: [OpenMP] Add more pass-through functions in DeviceTy.
Jul 24 2020, 6:57 AM · Restricted Project

Jul 23 2020

ye-luo created D84487: [OpenMP] Add more pass-through functions in DeviceTy.
Jul 23 2020, 11:02 PM · Restricted Project

Jul 22 2020

ye-luo accepted D84381: [OpenMP] Wait for kernel prior to memory deallocation.
Jul 22 2020, 9:48 PM · Restricted Project
ye-luo added a comment to D84381: [OpenMP] Wait for kernel prior to memory deallocation.

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.

Jul 22 2020, 9:46 PM · Restricted Project
ye-luo added a comment to D84381: [OpenMP] Wait for kernel prior to memory deallocation.

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.

Jul 22 2020, 8:12 PM · Restricted Project
ye-luo added a comment to D84381: [OpenMP] Wait for kernel prior to memory deallocation.

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?

Jul 22 2020, 7:59 PM · Restricted Project

Jul 21 2020

ye-luo updated subscribers of D84182: [OPENMP]Fix PR46012: declare target pointer cannot be accessed in target region..
Jul 21 2020, 8:30 AM · Restricted Project
ye-luo accepted D84182: [OPENMP]Fix PR46012: declare target pointer cannot be accessed in target region..

I verified that 46012 is fixed with this patch

Jul 21 2020, 8:29 AM · Restricted Project

Jul 19 2020

ye-luo added a comment to D82719: [OpenMPOpt][SplitMemTransfer][WIP] Getting values stored in offload arrays.

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?

Jul 19 2020, 6:54 PM · Restricted Project, Restricted Project

Jul 16 2020

ye-luo added a comment to D83963: [OpenMP] Use `abort` not `error` for fatal runtime exceptions.

Add flush

Jul 16 2020, 12:58 PM · Restricted Project
ye-luo added a comment to D83963: [OpenMP] Use `abort` not `error` for fatal runtime exceptions.

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

Jul 16 2020, 11:48 AM · Restricted Project
ye-luo added a comment to D83963: [OpenMP] Use `abort` not `error` for fatal runtime exceptions.

Are you sure abort gets fprintf flushed?

Jul 16 2020, 11:01 AM · Restricted Project

Jul 14 2020

ye-luo accepted D83707: [OpenMP][NFC] Emit remarks during GPU state machine optimization.

LGTM

Jul 14 2020, 1:36 PM · Restricted Project, Restricted Project

Jul 13 2020

ye-luo added a comment to D83718: [CallGraph] Update callback call sites in RefreshCallGraph.

Assertion failure in PR46687 goes away with this patch.

Jul 13 2020, 4:25 PM · Restricted Project
ye-luo added inline comments to D83707: [OpenMP][NFC] Emit remarks during GPU state machine optimization.
Jul 13 2020, 12:08 PM · Restricted Project, Restricted Project

Jul 7 2020

ye-luo added inline comments to D82718: [OpenMP] Use primary context in CUDA plugin.
Jul 7 2020, 6:13 AM · Restricted Project
ye-luo updated the diff for D82718: [OpenMP] Use primary context in CUDA plugin.
Jul 7 2020, 6:12 AM · Restricted Project

Jul 6 2020

ye-luo added inline comments to D82718: [OpenMP] Use primary context in CUDA plugin.
Jul 6 2020, 9:01 PM · Restricted Project

Jul 2 2020

ye-luo added inline comments to D82718: [OpenMP] Use primary context in CUDA plugin.
Jul 2 2020, 10:17 AM · Restricted Project
ye-luo updated the diff for D82718: [OpenMP] Use primary context in CUDA plugin.
Jul 2 2020, 10:14 AM · Restricted Project
ye-luo added a comment to D82718: [OpenMP] Use primary context in CUDA plugin.

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

Jul 2 2020, 6:25 AM · Restricted Project

Jul 1 2020

ye-luo added a comment to D82718: [OpenMP] Use primary context in CUDA plugin.
Jul 1 2020, 12:26 PM · Restricted Project

Jun 30 2020

ye-luo updated the diff for D82718: [OpenMP] Use primary context in CUDA plugin.

Updated description regarding setting CU_CTX_SCHED_BLOCKING_SYNC.

Jun 30 2020, 6:59 PM · Restricted Project
ye-luo added a comment to D82718: [OpenMP] Use primary context in CUDA plugin.

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.

Jun 30 2020, 1:35 PM · Restricted Project
ye-luo added a comment to D82718: [OpenMP] Use primary context in CUDA plugin.

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?

Jun 30 2020, 1:35 PM · Restricted Project
ye-luo added a comment to D82718: [OpenMP] Use primary context in CUDA plugin.

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

Jun 30 2020, 8:06 AM · Restricted Project

Jun 29 2020

ye-luo updated the diff for D82718: [OpenMP] Use primary context in CUDA plugin.
Jun 29 2020, 8:24 PM · Restricted Project
ye-luo added a comment to D82718: [OpenMP] Use primary context in CUDA plugin.

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

Jun 29 2020, 4:55 PM · Restricted Project
ye-luo added a comment to D82789: [OpenMP] fix clang warning about printf format in CUDA plugin.

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

Jun 29 2020, 1:36 PM · Restricted Project
ye-luo created D82789: [OpenMP] fix clang warning about printf format in CUDA plugin.
Jun 29 2020, 10:48 AM · Restricted Project
ye-luo added a comment to D82718: [OpenMP] Use primary context in CUDA plugin.

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.

Jun 29 2020, 9:42 AM · Restricted Project
ye-luo added a comment to D82718: [OpenMP] Use primary context in CUDA plugin.

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

Jun 29 2020, 9:42 AM · Restricted Project