Page MenuHomePhabricator

grokos (George Rokos)
User

Projects

User does not belong to any projects.

User Details

User Since
Feb 16 2016, 1:22 PM (173 w, 5 d)

Recent Activity

Fri, Jun 14

grokos accepted D63106: [OpenMP][libomptarget] Add support for declare target to clause under unified memory.

LGTM

Fri, Jun 14, 12:43 PM · Restricted Project
grokos added inline comments to D63106: [OpenMP][libomptarget] Add support for declare target to clause under unified memory.
Fri, Jun 14, 11:18 AM · Restricted Project

Tue, Jun 4

grokos added a comment to D62841: [openmp] Use libffi only when LLVM_ENABLE_FFI is on.

May I propose something extra? I guess the purpose of this patch is to disable libffi, for instance on systems where the library is not installed or for whatever reason the user doesn't want to use it. In that case, we should accommodate for the standalone build as well. We could introduce a new Cmake variable LIBOMPTARGET_ENABLE_FFI and

  • set it to LLVM_ENABLE_FFI if building in-tree or
  • have a default value of true if building standalone and let the user set it to false if they want to

Any thoughts?

Tue, Jun 4, 3:06 PM · Restricted Project
grokos added a comment to D60972: [OpenMP 5.0] libomptarget interface for declare mapper functions.

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.

Sure, I do prefer __tgt, as long as everyone else is okay with it.

Tue, Jun 4, 12:00 PM · Restricted Project

Mon, Jun 3

grokos accepted D62393: [OPENMP][NVPTX]Mark parallel level counter as volatile..

I formatted the description a bit. The patch looks good and the reasoning behind it is now obvious, in line with Nvidia's documentation.

Mon, Jun 3, 5:29 PM · Restricted Project
grokos updated the summary of D62393: [OPENMP][NVPTX]Mark parallel level counter as volatile..
Mon, Jun 3, 5:28 PM · Restricted Project
grokos added a comment to D60972: [OpenMP 5.0] libomptarget interface for declare mapper functions.

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!

Mon, Jun 3, 1:54 PM · Restricted Project
grokos added inline comments to D62397: [OPENMP][NVPTX]Relax flush directive..
Mon, Jun 3, 1:45 PM · Restricted Project
grokos accepted D60884: [OpenMP][libomptarget] Enable usage of unified memory for declare target link variables.

OK, there don't seem to be any more comments, looks good.

Mon, Jun 3, 1:43 PM · Restricted Project, Restricted Project
grokos added inline comments to D60884: [OpenMP][libomptarget] Enable usage of unified memory for declare target link variables.
Mon, Jun 3, 11:31 AM · Restricted Project, Restricted Project

Fri, May 31

grokos added inline comments to D62397: [OPENMP][NVPTX]Relax flush directive..
Fri, May 31, 3:22 PM · Restricted Project
grokos added inline comments to D62397: [OPENMP][NVPTX]Relax flush directive..
Fri, May 31, 3:13 PM · Restricted Project
grokos added inline comments to D60884: [OpenMP][libomptarget] Enable usage of unified memory for declare target link variables.
Fri, May 31, 3:00 PM · Restricted Project, Restricted Project

Wed, May 29

grokos added a comment to D56274: [OPENMP][NVPTX]Fix incompatibility of __syncthreads with LLVM, NFC..
Wed, May 29, 12:10 PM · Restricted Project
grokos added a comment to D62581: CallSiteSplitting: Respect convergent.

Also, can you include a reference in the commit message pointing to D56274 as this was the case demonstrating the need for this fix? Thanks!

Wed, May 29, 12:09 PM

Sat, May 25

grokos added inline comments to D60884: [OpenMP][libomptarget] Enable usage of unified memory for declare target link variables.
Sat, May 25, 12:45 PM · Restricted Project, Restricted Project

Fri, May 24

grokos added a comment to D62199: [OPENMP][NVPTX]Fix barriers and parallel level counters, NFC..

I agree with some of your complaints, I disagree with others:

Fri, May 24, 2:56 AM · Restricted Project, Restricted Project

Thu, May 23

grokos requested changes to D62199: [OPENMP][NVPTX]Fix barriers and parallel level counters, NFC..

Let me address what has been written so far from my perspective and request some changes since there has been demand for them:

  1. No doubt there are 2 changes here, unrelated to one another. Yes, they could be split into 2 patches to strictly comply with community guidelines. And I would prefer it if the author had submitted two distinct patches. On the other hand, I don't like missing the forest for the trees. This is a 4-line patch, I figured out it is OK to let both changes in at once. Sure, the small size of a patch is not a free pass to doing whatever we want, but in this case committing both changes in one go doesn't do any harm, at least none that I can think of so that I would make a big deal out of it. There's the letter of the law and the spirit of the law, in cases where these two collide I tend to go with the latter.
  2. My bad that I didn't request a test for cache flushing, I was predisposed due to the NFC tag. I agree there is no way to test whether the cache has been flushed, but @ABataev you could provide a test case which fails without cache flushing so that you can demonstrate that this patch does indeed work. There must be such a test case, after all how did you figure out there is a problem there?
  3. The NFC bit does not apply to the barrier fix. If we keep both changes in one patch, that should be reflected in the title. I agree that not being able to test a change does not make it NFC. You can call it NTC (non-testable change), but not NFC.
  4. In case of volatile, it seems there's a gray line here between what constitutes a NFC patch. The way I see it, "volatile" prevents compiler optimizations, it doesn't change anything in the way this code works, the algorithms/data structures used etc., the only change is how the code is compiled. On the other hand, we could argue that since the compiled library behaves differently there is indeed a functional change. Once again, to me this is a gray line and there is no absolute answer. However, since here there are more objections than approvals, @ABataev can you change the patch's title and remove the NFC part altogether?
  5. Regarding justification for the need for volatile, I agree that the explanation of the problem is very short, but the solution makes sense given the nature of that problem - at least to me. I could have blocked the patch until the author finds a way to describe the issue in detail - in the meantime there would be a correctness problem in the library which could cause inexplicable failures/wrong results to the users. Even if the "fix" is more of a hack - which I don't believe - I would rather let this hack in than leave the correctness problem lying around indefinitely. From the description I understand that a thread may cache values from the parallelLevel array but at the same time those values change in RAM without the thread knowing about it. This leads to corrupt states where some threads have a stale view of the parallel level. The only way to demonstrate the problem is to post ptx code, which I think is too much in the scope of a review, that's why I don't see the point in providing more details. The author could point to two places in the code, one where a thread caches a value from that array and another one where another thread modifies the global value, but it is not always possible to track where those changes occur. This is in case I understood the problem correctly. If not, then @ABataev I am also requesting a detailed explanation of the problem because the current description leads to wrong conclusions.
  6. 1, 4 and 5 represent my point of view, which is of course subjective, anyone could have a different mentality. But calling my reviews "reviews" and my activity here "suspicious" doesn't contribute at all. In pretty much the same way the "this is your problem, I've got the owner's approval" reply is inappropriate (and in no way do I endorse this "argument" as a reply to reviewers' requests).
  7. Since February 2018 I have been the owner of the libomptarget project. The issue of libomptarget having no owner came up in one of our biweekly multi-company telecons, I was the only one who volunteered and at a subsequent F2F there was consensus for my appointment since I had written most of the code for the base library and the CUDA and generic-elf64 plugins.
  8. Regarding the nvptx device runtime, certainly it is a different project from the rest of libomptarget. I think the RTL is still considered "coupled" with libomptarget for historical reasons only (both the base library, the plugins and the RTL were all written by the same people at the same company in the scope of the same project). But in terms of further development and expertise needed, it is a completely different kettle of fish. From my current position I can only contribute to the RTL by reviewing patches, it is true that I no longer do any development in this RTL. If anyone feels that the nvptx RTL repository should be managed/owned separately from the rest of libomptarget, please speak up and I will raise the issue at the next community telecon.
Thu, May 23, 8:55 PM · Restricted Project, Restricted Project

Tue, May 21

grokos accepted D62199: [OPENMP][NVPTX]Fix barriers and parallel level counters, NFC..

LGTM

Tue, May 21, 11:29 AM · Restricted Project, Restricted Project

Mon, May 20

grokos accepted D60223: [OpenMP][libomptarget] Enable requires flags for target libraries..

I don't have any other comments or suggestions, the patch looks good.

Mon, May 20, 12:29 PM · Restricted Project

May 17 2019

grokos added inline comments to D60223: [OpenMP][libomptarget] Enable requires flags for target libraries..
May 17 2019, 12:56 PM · Restricted Project

May 10 2019

grokos accepted D61801: [OPENMP][NVPTX]Simplify handling of thread limit, NFC..

LGTM

May 10 2019, 7:10 PM · Restricted Project
grokos accepted D61785: [OPENMP][NVPTX]Improve number of threads counter, NFC..

LGTM

May 10 2019, 11:49 AM · Restricted Project, Restricted Project

May 3 2019

grokos accepted D61526: [OPENMP][NVPTX]Improve thread limit counter, NFC..

Looks good.

May 3 2019, 12:42 PM · Restricted Project, Restricted Project
grokos accepted D61523: [OPENMP][NVPTX]Improve number of threads counter, NFC..

Looks good.

May 3 2019, 12:42 PM · Restricted Project

May 2 2019

grokos accepted D61459: [OPENMP][NVPTX]Improved several standard OpenMP functions, NFC..

Looks good.

May 2 2019, 4:35 PM · Restricted Project, Restricted Project
grokos accepted D61395: [OPENMP][NVPTX]Improve code by using parallel level counter..

LGTM

May 2 2019, 12:52 PM · Restricted Project, Restricted Project

Apr 26 2019

grokos accepted D60918: [OPENMP][NVPTX]Correctly handle L2 parallelism in SPMD mode..

I think it looks good now.

Apr 26 2019, 12:20 PM · Restricted Project, Restricted Project

Apr 24 2019

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

I see that target_data_mapper pretty much copies the entire code from target_data_begin and target_data_end. That's 200+ duplicate lines of code. Couldn't we just add this check:

Apr 24 2019, 12:12 PM · Restricted Project
grokos added a reviewer for D60972: [OpenMP 5.0] libomptarget interface for declare mapper functions: grokos.
Apr 24 2019, 12:00 PM · Restricted Project

Mar 25 2019

grokos accepted D55952: [libomptarget] Introduce LIBOMPTARGET_ENABLE_DEBUG cmake option..

OK, now it makes sense. I had forgotten what was in the description since the patch was first submitted. :)

Mar 25 2019, 7:38 PM · Restricted Project, Restricted Project

Mar 22 2019

grokos added a comment to D55952: [libomptarget] Introduce LIBOMPTARGET_ENABLE_DEBUG cmake option..

This revision looks much better. Please have a look at the inline comment.

Mar 22 2019, 8:08 PM · Restricted Project, Restricted Project

Jan 4 2019

grokos accepted D56332: [OPENMP][NVPTX]Fix dynamic scheduling..

Nice approach to optimizing the degree of collisions of atomic adds. LGTM.

Jan 4 2019, 7:48 PM
grokos accepted D56290: [OPENMP][NVPTX]General formatting/code improvement, NFC..

Looks good. Thanks!

Jan 4 2019, 12:01 PM
grokos accepted D56278: [OPENMP][NVPTX]Improve performance + reduce number of used registers..

LGTM. Thanks for all this work!

Jan 4 2019, 8:25 AM
grokos added inline comments to D56290: [OPENMP][NVPTX]General formatting/code improvement, NFC..
Jan 4 2019, 8:24 AM
grokos added inline comments to D56278: [OPENMP][NVPTX]Improve performance + reduce number of used registers..
Jan 4 2019, 8:06 AM

Jan 3 2019

grokos added inline comments to D55588: [OpenMP] Fix nvidia-cuda-toolkit detection on Debian/Ubuntu.
Jan 3 2019, 5:44 PM
grokos accepted D56274: [OPENMP][NVPTX]Fix incompatibility of __syncthreads with LLVM, NFC..

I'll accept the patch for the sake of consistency and correctness of execution. Just one question:

which cost is too high

So should we expect a performance penalty until function copy is fixed in LLVM and we can revert back to __syncthreads()?

Jan 3 2019, 8:38 AM · Restricted Project
grokos accepted D55588: [OpenMP] Fix nvidia-cuda-toolkit detection on Debian/Ubuntu.

LGTM.

Jan 3 2019, 5:12 AM

Dec 31 2018

grokos added a comment to D55952: [libomptarget] Introduce LIBOMPTARGET_ENABLE_DEBUG cmake option..

I'm not sure what we are trying to do here. OMPTARGET_DEBUG is not an environment variable, it's a preprocessor definition. It is used by libomptarget/include/omptarget.h to include all necessary header files and define the DEBUGP macro appropriately, which in turn is used by libomptarget/src/private.h to define the DP macro used throughout the library.

Dec 31 2018, 9:30 AM · Restricted Project, Restricted Project

Dec 27 2018

grokos accepted D56100: [OPENMP][NVPTX]Fixed initialization of the data-sharing interface..
Dec 27 2018, 4:21 PM
grokos accepted D56102: [OPENMP][NVPTX]Added/fixed debugging messages, NFC..

LGTM.

Dec 27 2018, 8:26 AM
grokos added inline comments to D56102: [OPENMP][NVPTX]Added/fixed debugging messages, NFC..
Dec 27 2018, 8:20 AM
grokos accepted D56101: [OPENMP][NVPTX]Outline assert into noinline function, NFC..

Looks good.

Dec 27 2018, 8:19 AM
grokos added inline comments to D56100: [OPENMP][NVPTX]Fixed initialization of the data-sharing interface..
Dec 27 2018, 8:11 AM

Dec 18 2018

grokos added a comment to D55588: [OpenMP] Fix nvidia-cuda-toolkit detection on Debian/Ubuntu.

Since a similar solution was adopted for clang, I think we should let this one land. After all, it's a matter of consistency between the two projects.

Dec 18 2018, 7:16 PM

Nov 5 2018

grokos edited reviewers for D54079: [OPENMP] Add Hurd support, added: AndreyChurbanov, tlwilmar, jlpeyton, hbae; removed: gtbercea, caomhin, grokos.
Nov 5 2018, 6:47 AM

Oct 22 2018

grokos added a comment to D46185: [OpenMP] Allow nvptx sm_30 to be used as an offloading device.

Are we still interested in this patch or should we abandon it?

Oct 22 2018, 6:28 AM · Restricted Project

Oct 11 2018

grokos accepted D53141: [OpenMP][libomptarget] Add runtime function for pushing coalesced global records.

Looks good.

Oct 11 2018, 8:41 AM

Oct 1 2018

grokos accepted D52725: [libomptarget-nvptx] reduction: Determine if runtime uninitialized.

Looks good. Thanks for taking care of this issue!

Oct 1 2018, 7:16 AM
grokos accepted D52701: [libomptarget-nvptx] Enable asserts in bclib.

Looks good. I have no idea why this option was left there...

Oct 1 2018, 1:44 AM

Sep 29 2018

grokos accepted D51787: [libomptarget-nvptx] Fix ancestor_thread_num and team_size (non-SPMD).

Looks good.

Sep 29 2018, 12:35 PM

Sep 28 2018

grokos accepted D52655: [libomptarget-nvptx] Align data sharing stack.

Interesting. Are there other cases where alignment is needed?

Sep 28 2018, 9:18 AM
grokos accepted D51785: [libomptarget-nvptx] Ignore calls to dynamic API.

LG

Sep 28 2018, 8:39 AM
grokos accepted D51783: [libomptarget-nvptx] Fix number of threads in parallel.

LG

Sep 28 2018, 8:39 AM
grokos accepted D51687: [libomptarget-nvptx] Add testing infrastructure.

Then this revision is good to go!

Sep 28 2018, 6:38 AM

Sep 21 2018

grokos added a comment to D51875: [OPENMP][NVPTX] Add support for lastprivates/reductions handling in SPMD constructs with lightweight runtime..

@Hahnfeld: Are the latest changes in line with your requirements/plans to reduce the memory footprint of the nvptx runtime?

Sep 21 2018, 7:14 AM

Sep 12 2018

grokos accepted D50188: [OpenMP][libomptarget] Simplify warp master selection for data sharing.

I completely missed this one. The proposed change makes sense.

Sep 12 2018, 6:23 AM

Sep 11 2018

grokos added inline comments to D51875: [OPENMP][NVPTX] Add support for lastprivates/reductions handling in SPMD constructs with lightweight runtime..
Sep 11 2018, 7:53 AM

Sep 10 2018

grokos added a comment to D51687: [libomptarget-nvptx] Add testing infrastructure.

Great job, adding testing infrastructure has been on our list for a long time!

Sep 10 2018, 4:22 PM

Sep 6 2018

grokos added a comment to D51107: [LIBOMPTARGET] Add support for mapping of lambda captures..

BTW, the main problem here is that there is no mapping between host pointer and device pointer for the firstprivates, just like for the mapped data. We just need to rework the way we allocate the memory for the firstprivate to make MEMBER_OF to work.

I don't see this path taken in the latest patch. Is there a major drawback of doing this?

I did it using new tgtArgsPositions, which handles mappings between the firstprivates and the corresponding target pointers.

I think it's more consistent to handle all data mapping in a single loop (in target_data_begin) as discussed before.

I think this is a different problem that should be solved separately. It requires a redesign of the existing solution, while I'm just trying to implement a new feature.

Sep 6 2018, 8:51 AM

Sep 4 2018

grokos accepted D51624: [libomptarget][CUDA] Use cuDeviceGetAttribute, NFCI..

Looks good.

Sep 4 2018, 8:08 AM
grokos accepted D51623: [libomptarget] PR38704: Fix erase of ShadowPtrMap.

Good catch, thanks for the fix!

Sep 4 2018, 8:00 AM
grokos added inline comments to D51622: [libomptarget][NVPTX] Drop dead code and data structures, NFCI..
Sep 4 2018, 7:54 AM

Sep 3 2018

grokos added a comment to D51107: [LIBOMPTARGET] Add support for mapping of lambda captures..

@AlexEichenberger: Have you confirmed whether flag 0x400 is free to be used for LAMBDA_REF or has been reserved for something else in XL?

Sep 3 2018, 2:52 PM

Aug 23 2018

grokos added a comment to D51107: [LIBOMPTARGET] Add support for mapping of lambda captures..

Is this patch about lambdas inside target regions or about lambdas containing a target region? From the description I assume it's the former. In that case, what happens if there are more than 1 lambdas in the target region?

  1. This is for lambdas inside target region.
  2. For each lambda we generate this sequence: <lambda1>, <capture1_1>, .., <capture1_n>, <lambda2>, <capture2_1>, ..., <capture2_m>, ...
Aug 23 2018, 5:53 AM

Aug 22 2018

grokos added a comment to D51107: [LIBOMPTARGET] Add support for mapping of lambda captures..

Is this patch about lambdas inside target regions or about lambdas containing a target region? From the description I assume it's the former. In that case, what happens if there are more than 1 lambdas in the target region?

Aug 22 2018, 11:43 AM

Aug 16 2018

grokos abandoned D44522: [Libomptarget] Full implementation of the target-offload-icv.
Aug 16 2018, 9:10 AM · Restricted Project

Aug 2 2018

grokos added a comment to D50158: [OpenMP] Add placeholder functions for the depend and nowait depend clauses for target data directives..

These functions are 1:1 the same as the current _nowait functions. If that's for handling depend clauses without nowait I think the compiler can just add another runtime call into libomp.

Aug 2 2018, 9:03 AM

Jul 20 2018

grokos added inline comments to D49564: [OPNEMP, NVPTX] Fixed sychronization construct + code cleanup..
Jul 20 2018, 4:54 PM

Jul 19 2018

grokos committed rOMP337455: [OpenMP][libomptarget] New map interface: remove translation code and ensure….
[OpenMP][libomptarget] New map interface: remove translation code and ensure…
Jul 19 2018, 6:48 AM
grokos committed rL337455: [OpenMP][libomptarget] New map interface: remove translation code and ensure….
[OpenMP][libomptarget] New map interface: remove translation code and ensure…
Jul 19 2018, 6:46 AM
grokos closed D44186: [OpenMP][libomptarget] New map interface: remove translation code and ensure proper alignment of struct members.
Jul 19 2018, 6:46 AM · Restricted Project
grokos added a comment to D44186: [OpenMP][libomptarget] New map interface: remove translation code and ensure proper alignment of struct members.

@ABataev: Can you put a link to the clang-side patch in the description so that we link the two patches together? Also, please let me know when you commit the clang patch so that I commit this one as well.

Jul 19 2018, 6:35 AM · Restricted Project
grokos added inline comments to D44186: [OpenMP][libomptarget] New map interface: remove translation code and ensure proper alignment of struct members.
Jul 19 2018, 6:31 AM · Restricted Project
grokos updated the diff for D44186: [OpenMP][libomptarget] New map interface: remove translation code and ensure proper alignment of struct members.

Added example code to demonstrate the need for padding in partially mapped structs.

Jul 19 2018, 6:24 AM · Restricted Project
grokos retitled D44186: [OpenMP][libomptarget] New map interface: remove translation code and ensure proper alignment of struct members from [OpenMP] New clang/libomptarget map interface: remove translation code to [OpenMP][libomptarget] New map interface: remove translation code and ensure proper alignment of struct members.
Jul 19 2018, 5:48 AM · Restricted Project

Jul 18 2018

grokos added a comment to D44186: [OpenMP][libomptarget] New map interface: remove translation code and ensure proper alignment of struct members.

Hi George, when do you plan to commit this patch? I cleaned the patch for the clang and it is ready to be committed.

Jul 18 2018, 10:15 AM · Restricted Project

Jul 17 2018

grokos accepted D49418: Also support several images for elf.

LGTM

Jul 17 2018, 9:08 AM

Jul 13 2018

grokos accepted D49204: [OpenMP][libomptarget] Fix data sharing and globalization infrastructure to work in SPMD mode.

I don't have any further comments.

Jul 13 2018, 2:08 AM

Jul 12 2018

grokos added inline comments to D49204: [OpenMP][libomptarget] Fix data sharing and globalization infrastructure to work in SPMD mode.
Jul 12 2018, 9:17 AM

Jul 9 2018

grokos accepted D49036: [OPENMP, NVPTX] Support several images in the executable..

Looks good. Do you think we should do the same for the generic-elf-64 plugin?

Jul 9 2018, 10:41 AM

Jun 29 2018

grokos accepted D48732: [OPENMP, NVPTX] Sync threads before start ordered loops..

The addition looks good. I've seen this piece of code in quite a few places, maybe we need to refactor it into a mini-function. But let's revisit that in the future.

Jun 29 2018, 9:12 AM

Jun 25 2018

grokos accepted D48480: [OPENMP, NVPTX] Fixes for NVPTX RTL.

Looks good to me, thanks for submitting those fixes!

Jun 25 2018, 5:43 AM

Jun 18 2018

grokos accepted D48286: [OpenMP] [CUDA] Expose teamid to the debug path.
Jun 18 2018, 11:20 AM · Restricted Project
grokos added a comment to D48286: [OpenMP] [CUDA] Expose teamid to the debug path.

LGTM.

Jun 18 2018, 11:20 AM · Restricted Project

May 24 2018

grokos committed rOMP333225: [CUDA]Fix dynamic|guided scheduling..
[CUDA]Fix dynamic|guided scheduling.
May 24 2018, 2:53 PM
grokos committed rL333225: [CUDA]Fix dynamic|guided scheduling..
[CUDA]Fix dynamic|guided scheduling.
May 24 2018, 2:16 PM
grokos closed D47333: [CUDA]Fix dynamic|guided scheduling..
May 24 2018, 2:16 PM
grokos accepted D47333: [CUDA]Fix dynamic|guided scheduling..

I think this looks good. Definitely a correct implementation compared to what we have now!

May 24 2018, 10:23 AM

May 22 2018

grokos accepted D47131: [libomptarget-nvptx] loop: Determine if runtime uninitialized.

Looks good. I don't see any reason why performance would be affected and it is a necessary change for the sake of correctness.

May 22 2018, 8:13 AM

May 21 2018

grokos accepted D47130: [CMake] Unify install path for libraries.

LGTM.

May 21 2018, 7:51 AM
grokos added a comment to D44186: [OpenMP][libomptarget] New map interface: remove translation code and ensure proper alignment of struct members.

Was the new map interface added to Clang and just missed it?

May 21 2018, 7:46 AM · Restricted Project

May 16 2018

grokos accepted D46930: [libomptarget-nvptx-bc] Pass found CUDA installations.

LGTM

May 16 2018, 8:32 AM
grokos accepted D46901: [libomptarget-nvptx] Test bitcode compiler flags and enable by default.

Looks good! Thanks for taking care of it!

May 16 2018, 8:31 AM

May 15 2018

grokos accepted D46840: [OpenMP][libomptarget] Add function for checking SPMD mode.

Looks good.

May 15 2018, 8:13 AM

May 14 2018

grokos requested changes to D46842: [OpenMP][libomptarget] Make bitcode library building depend on clang and llvm-linker being available .

I agree with Jonas, the initial version of the nvptx RTL that we tried to upstream was somewhat like what you try to do here and the consensus was that we do not want to have these dependencies.

May 14 2018, 1:50 PM

May 4 2018

grokos retitled D46210: [OpenMP] Use LIBOMPTARGET_DEVICE_RTL_DEBUG env var to control debug messages on the device side from [OpenMP] Use LIBOMPDEVICE_DEBUG env var to control debug messages on the device side to [OpenMP] Use LIBOMPTARGET_DEVICE_RTL_DEBUG env var to control debug messages on the device side.
May 4 2018, 12:06 PM · Restricted Project
grokos accepted D46210: [OpenMP] Use LIBOMPTARGET_DEVICE_RTL_DEBUG env var to control debug messages on the device side.

LGTM

May 4 2018, 12:06 PM · Restricted Project