ABataev (Alexey Bataev)
User

Projects

User does not belong to any projects.

User Details

User Since
Jan 31 2013, 4:40 AM (289 w, 3 d)

Recent Activity

Yesterday

ABataev updated subscribers of D50939: [WIP][OpenMP] Propagate 'declare target' information to device.

D38798 does not work correctly and there is no stable Sema-based solution to do this absolutely correctly. We need to adapt the codegen so that it emits metadata for the implicit declare target functions.

Sat, Aug 18, 7:50 AM

Fri, Aug 17

ABataev added a comment to D50845: [CUDA/OpenMP] Define only some host macros during device compilation.

Right, warning wasn't a good thought. We really want strict checking and would have to error out when we find a function that wasn't implicitly declare target on the host.
I meant to ask how common that would be? If that's only some known functions we could handle them separately.

Again, it does not matter how common is this situation. We cannot rely on the probability here, we need to make the compiler to work correctly in all possible situation, no matter how often they can occur.

Got that, I agree on the conservative approach: If we find a function to be called that wasn't checked (because it wasn't implicitly declare target on the host) the compiler can error out. That should be correct in all cases, shouldn't it?

There's a trade-off here:

  • How many TUs pass full analysis and how many don't? (today's situation; we know that some headers don't work)
  • How many TUs pass when we only check called functions (and error if we call non-checked ones) and how many regress compared to today's situation? If the number of regressions is zero for all practical situations but we can compile some important cases, that should be a win.
Fri, Aug 17, 9:17 AM
ABataev added a comment to D50845: [CUDA/OpenMP] Define only some host macros during device compilation.

Right, warning wasn't a good thought. We really want strict checking and would have to error out when we find a function that wasn't implicitly declare target on the host.
I meant to ask how common that would be? If that's only some known functions we could handle them separately.

Fri, Aug 17, 8:05 AM
ABataev added a comment to D50845: [CUDA/OpenMP] Define only some host macros during device compilation.

So ideally I think Clang should determine which functions are really declare target (either explicit or implicit) and only run semantical analysis on them. If a function is then found to be "broken" it's perfectly desirable to error back to the user.

It is not possible for OpenMP because we support implicit declare target functions. Clang cannot identify whether the function is going to be used on the device or not during sema analysis.

You are right, we can't do this during device compilation because we don't have an AST before Sema.

However I'm currently thinking about the following:

  1. Identify explicit and implicit declare target functions during host Sema and CodeGen.
  2. Attach meta-data for all of them to LLVM IR .bc which is passed via -fopenmp-host-ir-file-path. I think we already do something similar for outlined target regions?
  3. During device Sema query that meta-data so Clang knows when a function will be called from within a target region. Skip analysis of functions that are not needed for the device, just as CUDA does.
  4. Check that we don't need functions that weren't marked in 2. That's to catch users doing something like: `lang=c #pragma omp target { #ifdef NVPTX target_func() #endif } `

    For now that's just an idea, I didn't start implementing any of this yet. Do you think that could work?

I thought about this approach already. But it won't work in general. The main problem here is that host and device compilation phases may end up with the different set of implicit declare target functions. The main problem here not the user code, but the system libraries, which may use the different set of functions.

How common is that for functions that are used in target regions? In the worst case we can make my fourth point a warning and lose Sema checking for those functions.

Fri, Aug 17, 6:32 AM
ABataev added a comment to D50845: [CUDA/OpenMP] Define only some host macros during device compilation.

So ideally I think Clang should determine which functions are really declare target (either explicit or implicit) and only run semantical analysis on them. If a function is then found to be "broken" it's perfectly desirable to error back to the user.

It is not possible for OpenMP because we support implicit declare target functions. Clang cannot identify whether the function is going to be used on the device or not during sema analysis.

You are right, we can't do this during device compilation because we don't have an AST before Sema.

However I'm currently thinking about the following:

  1. Identify explicit and implicit declare target functions during host Sema and CodeGen.
  2. Attach meta-data for all of them to LLVM IR .bc which is passed via -fopenmp-host-ir-file-path. I think we already do something similar for outlined target regions?
  3. During device Sema query that meta-data so Clang knows when a function will be called from within a target region. Skip analysis of functions that are not needed for the device, just as CUDA does.
  4. Check that we don't need functions that weren't marked in 2. That's to catch users doing something like: `lang=c #pragma omp target { #ifdef NVPTX target_func() #endif } `

    For now that's just an idea, I didn't start implementing any of this yet. Do you think that could work?
Fri, Aug 17, 6:15 AM

Thu, Aug 16

ABataev added a comment to D50845: [CUDA/OpenMP] Define only some host macros during device compilation.

If I understand it correctly, the root cause of this exercise is that we want to compile for GPU using plain C. CUDA avoids this issue by separating device and host code via target attributes and clang has few special cases to ignore inline assembly errors in the host code if we're compiling for device. For OpenMP there's no such separation, not in the system headers, at least.

Yes, that's one of the nice properties of CUDA (for the compiler). There used to be the same restriction for OpenMP where all functions used in target regions needed to be put in declare target. However that was relaxed in favor of implicitly marking all called functions in that TU to be declare target.
So ideally I think Clang should determine which functions are really declare target (either explicit or implicit) and only run semantical analysis on them. If a function is then found to be "broken" it's perfectly desirable to error back to the user.

Thu, Aug 16, 12:02 PM
ABataev added a comment to D50845: [CUDA/OpenMP] Define only some host macros during device compilation.

Maybe for device compilation we also should define __NO_MATH_INLINES and __NO_STRING_INLINES macros to disable inline assembly in glibc?

The problem is that __NO_MATH_INLINES doesn't even avoid all inline assembly from bits/mathinline.h :-( incidentally Clang already defines __NO_MATH_INLINES for x86 (due to an old bug which has been fixed long ago) - and on CentOS we still have problems as described in PR38464.

As a second thought: This might be valid for NVPTX, but I don't think it's a good idea for x86-like offloading targets - they might well profit from inline assembly code.

Thu, Aug 16, 8:17 AM
ABataev added a comment to D50845: [CUDA/OpenMP] Define only some host macros during device compilation.

Maybe for device compilation we also should define __NO_MATH_INLINES and __NO_STRING_INLINES macros to disable inline assembly in glibc?

Thu, Aug 16, 8:04 AM
ABataev added a comment to D50840: [InstCombine] Extend collectShuffleElements to support extract/zext/insert patterns.

It could be the kind of thing we should do in slp @ABataev what do you think?

Yes, looks like the opportunity for the SLP Vectorizer.

Thu, Aug 16, 7:45 AM

Wed, Aug 15

ABataev committed rL339805: [OPENMP] FIx processing of declare target variables..
[OPENMP] FIx processing of declare target variables.
Wed, Aug 15, 12:46 PM
ABataev committed rC339805: [OPENMP] FIx processing of declare target variables..
[OPENMP] FIx processing of declare target variables.
Wed, Aug 15, 12:46 PM
ABataev added a comment to D46021: [DEBUGINFO] Disable emission of the dwarf sections, but allow directives..
In D46021#1200884, @tra wrote:

Is there a way to control this from clang command line?

Wed, Aug 15, 10:01 AM

Tue, Aug 14

ABataev committed rL339704: [OPENMP] Fix processing of declare target construct..
[OPENMP] Fix processing of declare target construct.
Tue, Aug 14, 11:32 AM
ABataev committed rC339704: [OPENMP] Fix processing of declare target construct..
[OPENMP] Fix processing of declare target construct.
Tue, Aug 14, 11:32 AM

Mon, Aug 13

ABataev committed rC339603: [OPENMP] Fix emission of the loop doacross constructs..
[OPENMP] Fix emission of the loop doacross constructs.
Mon, Aug 13, 12:05 PM
ABataev committed rL339603: [OPENMP] Fix emission of the loop doacross constructs..
[OPENMP] Fix emission of the loop doacross constructs.
Mon, Aug 13, 12:05 PM
ABataev committed rL339574: Revert "[OPENMP] Fix emission of the loop doacross constructs.".
Revert "[OPENMP] Fix emission of the loop doacross constructs."
Mon, Aug 13, 7:43 AM
ABataev committed rC339574: Revert "[OPENMP] Fix emission of the loop doacross constructs.".
Revert "[OPENMP] Fix emission of the loop doacross constructs."
Mon, Aug 13, 7:43 AM
ABataev committed rC339568: [OPENMP] Fix emission of the loop doacross constructs..
[OPENMP] Fix emission of the loop doacross constructs.
Mon, Aug 13, 7:06 AM
ABataev committed rL339568: [OPENMP] Fix emission of the loop doacross constructs..
[OPENMP] Fix emission of the loop doacross constructs.
Mon, Aug 13, 7:06 AM

Tue, Aug 7

ABataev committed rL339166: [SLP] Fix insert point for reused extract instructions..
[SLP] Fix insert point for reused extract instructions.
Tue, Aug 7, 12:21 PM
ABataev closed D50143: [SLP] Fix insert point for reused extract instructions..
Tue, Aug 7, 12:21 PM
ABataev updated the diff for D50143: [SLP] Fix insert point for reused extract instructions..

Added an extra test.

Tue, Aug 7, 9:53 AM
ABataev added inline comments to D50143: [SLP] Fix insert point for reused extract instructions..
Tue, Aug 7, 9:51 AM
ABataev committed rC339152: [OPENMP] Mark variables captured in declare target region as implicitly.
[OPENMP] Mark variables captured in declare target region as implicitly
Tue, Aug 7, 9:15 AM
ABataev committed rL339152: [OPENMP] Mark variables captured in declare target region as implicitly.
[OPENMP] Mark variables captured in declare target region as implicitly
Tue, Aug 7, 9:15 AM

Mon, Aug 6

ABataev added a comment to D50181: [SLP] Add bundle reordering function for already scheduled bundles. .

It really requires a test.

Mon, Aug 6, 7:38 AM

Fri, Aug 3

ABataev accepted D50218: [OpenMP] Encode offload target triples into comdat key for offload initialization code.

LG

Fri, Aug 3, 1:12 PM
ABataev added inline comments to D50218: [OpenMP] Encode offload target triples into comdat key for offload initialization code.
Fri, Aug 3, 11:28 AM
ABataev added inline comments to D50218: [OpenMP] Encode offload target triples into comdat key for offload initialization code.
Fri, Aug 3, 10:37 AM
ABataev updated the diff for D45784: [DEBUG_INFO, NVPTX] Fix relocation info..

Reworked the patch to change the way we get file id to make it more robust.

Fri, Aug 3, 7:34 AM
ABataev removed a reviewer for D48100: Append new attributes to the end of an AttributeList.: ABataev.
Fri, Aug 3, 7:03 AM

Thu, Aug 2

ABataev updated the diff for D45784: [DEBUG_INFO, NVPTX] Fix relocation info..

Reworked according to Eric's comments

Thu, Aug 2, 1:41 PM
ABataev added a comment to D50158: [OpenMP] Add placeholder functions for the depend and nowait depend clauses for target data directives..

We should not be having calls to _-kmpc_omp_taskwait in libomptarget.
The wait should be done in the user code and then libomptarget routine should be called.

There is no way to implement asynchronous offloading without having knowledge of the dependencies. Intel failed to keep the separation strict for liboffload, so I don't think it helps to complain about the situation.

Thu, Aug 2, 9:45 AM
ABataev updated the diff for D45822: [DEBUGINFO, NVPTX] Try to pack bytes data into a single string..

Reworked according to Eric's comments.

Thu, Aug 2, 7:18 AM

Wed, Aug 1

ABataev updated the diff for D46061: [DEBUGINFO, NVPTX] Disable emission of ',debug' option if only debug directives are allowed..

Updated and reworked to latest changes

Wed, Aug 1, 1:33 PM
ABataev committed rL338616: [DEBUGINFO] Disable emission of the dwarf sections, but allow directives..
[DEBUGINFO] Disable emission of the dwarf sections, but allow directives.
Wed, Aug 1, 12:39 PM
ABataev closed D46021: [DEBUGINFO] Disable emission of the dwarf sections, but allow directives..
Wed, Aug 1, 12:38 PM
ABataev updated the diff for D46021: [DEBUGINFO] Disable emission of the dwarf sections, but allow directives..

Address David's comments

Wed, Aug 1, 12:03 PM
ABataev added inline comments to D46021: [DEBUGINFO] Disable emission of the dwarf sections, but allow directives..
Wed, Aug 1, 11:50 AM
ABataev updated the diff for D46021: [DEBUGINFO] Disable emission of the dwarf sections, but allow directives..

Reworked to control output by DICompileUnit emission kind.

Wed, Aug 1, 10:42 AM
ABataev created D50143: [SLP] Fix insert point for reused extract instructions..
Wed, Aug 1, 7:00 AM

Tue, Jul 31

ABataev added inline comments to D49928: [SLP] Fix PR38339: Instruction does not dominate all uses!.
Tue, Jul 31, 1:57 PM
ABataev added inline comments to D49928: [SLP] Fix PR38339: Instruction does not dominate all uses!.
Tue, Jul 31, 1:25 PM
ABataev committed rL338413: [OPENMP] Change linkage of offloading symbols to support dropping.
[OPENMP] Change linkage of offloading symbols to support dropping
Tue, Jul 31, 11:28 AM
ABataev committed rC338413: [OPENMP] Change linkage of offloading symbols to support dropping.
[OPENMP] Change linkage of offloading symbols to support dropping
Tue, Jul 31, 11:28 AM
ABataev committed rL338399: [OPENMP] Prevent problems with linking of the static variables..
[OPENMP] Prevent problems with linking of the static variables.
Tue, Jul 31, 9:40 AM
ABataev committed rC338399: [OPENMP] Prevent problems with linking of the static variables..
[OPENMP] Prevent problems with linking of the static variables.
Tue, Jul 31, 9:40 AM
ABataev committed rL338380: [SLP] Fix PR38339: Instruction does not dominate all uses!.
[SLP] Fix PR38339: Instruction does not dominate all uses!
Tue, Jul 31, 7:03 AM
ABataev closed D49928: [SLP] Fix PR38339: Instruction does not dominate all uses!.
Tue, Jul 31, 7:03 AM

Mon, Jul 30

ABataev added a comment to D45784: [DEBUG_INFO, NVPTX] Fix relocation info..

nvcc always emits debug location before the very first label that represents function start. So, we can consider this as the requirement for NVPTX.

Mon, Jul 30, 3:21 PM
ABataev added a comment to D45784: [DEBUG_INFO, NVPTX] Fix relocation info..

But you have other locations inside of functions so I don't get it.

eric

Mon, Jul 30, 3:18 PM
ABataev added a comment to D45784: [DEBUG_INFO, NVPTX] Fix relocation info..

Can you explain this a bit more with some examples of wrong relocations?

Thanks!

-eric

Mon, Jul 30, 3:12 PM
ABataev committed rL338252: [OPENMP] Modify the info about OpenMP support in UsersManual, NFC..
[OPENMP] Modify the info about OpenMP support in UsersManual, NFC.
Mon, Jul 30, 7:44 AM
ABataev committed rC338252: [OPENMP] Modify the info about OpenMP support in UsersManual, NFC..
[OPENMP] Modify the info about OpenMP support in UsersManual, NFC.
Mon, Jul 30, 7:44 AM

Fri, Jul 27

ABataev committed rL338158: [DEBUG_INFO] Fix tests, NFC..
[DEBUG_INFO] Fix tests, NFC.
Fri, Jul 27, 1:18 PM
ABataev committed rC338158: [DEBUG_INFO] Fix tests, NFC..
[DEBUG_INFO] Fix tests, NFC.
Fri, Jul 27, 1:18 PM
ABataev committed rL338155: [DEBUGINFO] Disable unsupported debug info options for NVPTX target..
[DEBUGINFO] Disable unsupported debug info options for NVPTX target.
Fri, Jul 27, 12:51 PM
ABataev committed rC338155: [DEBUGINFO] Disable unsupported debug info options for NVPTX target..
[DEBUGINFO] Disable unsupported debug info options for NVPTX target.
Fri, Jul 27, 12:46 PM
This revision was not accepted when it landed; it landed in state Needs Review.
Fri, Jul 27, 12:46 PM
ABataev added a comment to D49148: [DEBUGINFO] Disable unsupported debug info options for NVPTX target..

Eric accepted the patch offline.

Fri, Jul 27, 12:18 PM
ABataev created D49928: [SLP] Fix PR38339: Instruction does not dominate all uses!.
Fri, Jul 27, 12:09 PM
ABataev committed rL338139: [OPENMP] Static variables on device must be externally visible..
[OPENMP] Static variables on device must be externally visible.
Fri, Jul 27, 10:37 AM
ABataev committed rC338139: [OPENMP] Static variables on device must be externally visible..
[OPENMP] Static variables on device must be externally visible.
Fri, Jul 27, 10:37 AM
ABataev updated the diff for D46021: [DEBUGINFO] Disable emission of the dwarf sections, but allow directives..

Updated to latest version

Fri, Jul 27, 8:11 AM
ABataev updated the diff for D45822: [DEBUGINFO, NVPTX] Try to pack bytes data into a single string..

Added maximal length of the merged directives.

Fri, Jul 27, 7:51 AM
ABataev updated the diff for D45784: [DEBUG_INFO, NVPTX] Fix relocation info..

Updated to the latest revision.

Fri, Jul 27, 6:49 AM

Thu, Jul 26

ABataev added inline comments to D49148: [DEBUGINFO] Disable unsupported debug info options for NVPTX target..
Thu, Jul 26, 1:30 PM
ABataev updated the diff for D49148: [DEBUGINFO] Disable unsupported debug info options for NVPTX target..

Address ERic's comments.

Thu, Jul 26, 1:29 PM
ABataev committed rC338055: [OPENMP, DOCS] Fixed typo, NFC..
[OPENMP, DOCS] Fixed typo, NFC.
Thu, Jul 26, 11:41 AM
ABataev committed rL338055: [OPENMP, DOCS] Fixed typo, NFC..
[OPENMP, DOCS] Fixed typo, NFC.
Thu, Jul 26, 11:41 AM
ABataev committed rL338049: [OPENMP] What's new for OpenMP in clang..
[OPENMP] What's new for OpenMP in clang.
Thu, Jul 26, 10:54 AM
ABataev committed rC338049: [OPENMP] What's new for OpenMP in clang..
[OPENMP] What's new for OpenMP in clang.
Thu, Jul 26, 10:54 AM
ABataev committed rL338039: [DEBUGINFO, NVPTX] Emit correct debug information for local variables..
[DEBUGINFO, NVPTX] Emit correct debug information for local variables.
Thu, Jul 26, 9:30 AM
ABataev closed D45963: [DEBUGINFO, NVPTX] Emit correct debug information for local variables..
Thu, Jul 26, 9:30 AM
ABataev committed rL338036: [DEBUGINFO, NVPTX] Set `DW_AT_frame_base` to `DW_OP_call_frame_cfa`..
[DEBUGINFO, NVPTX] Set `DW_AT_frame_base` to `DW_OP_call_frame_cfa`.
Thu, Jul 26, 9:10 AM
ABataev closed D45785: [DEBUGINFO, NVPTX] Set `DW_AT_frame_base` to `DW_OP_call_frame_cfa`..
Thu, Jul 26, 9:10 AM
ABataev committed rC338032: [OPENMP] Force OpenMP 4.5 when compiling for offloading..
[OPENMP] Force OpenMP 4.5 when compiling for offloading.
Thu, Jul 26, 8:18 AM
ABataev committed rL338032: [OPENMP] Force OpenMP 4.5 when compiling for offloading..
[OPENMP] Force OpenMP 4.5 when compiling for offloading.
Thu, Jul 26, 8:17 AM

Wed, Jul 25

ABataev committed rC337957: [OPENMP] ThreadId in serialized parallel regions is 0..
[OPENMP] ThreadId in serialized parallel regions is 0.
Wed, Jul 25, 1:03 PM
ABataev committed rL337957: [OPENMP] ThreadId in serialized parallel regions is 0..
[OPENMP] ThreadId in serialized parallel regions is 0.
Wed, Jul 25, 1:03 PM
ABataev committed rL337941: [OPENMP] Exclude service expressions/statements from the list of.
[OPENMP] Exclude service expressions/statements from the list of
Wed, Jul 25, 10:28 AM
ABataev committed rC337941: [OPENMP] Exclude service expressions/statements from the list of.
[OPENMP] Exclude service expressions/statements from the list of
Wed, Jul 25, 10:27 AM
ABataev committed rC337928: [OPENMP] Fix PR38256: Fix locations of the artificial conditional op..
[OPENMP] Fix PR38256: Fix locations of the artificial conditional op.
Wed, Jul 25, 7:40 AM
ABataev committed rL337928: [OPENMP] Fix PR38256: Fix locations of the artificial conditional op..
[OPENMP] Fix PR38256: Fix locations of the artificial conditional op.
Wed, Jul 25, 7:40 AM

Mon, Jul 23

ABataev committed rOMP337691: [OPNEMP, NVPTX] Fixed sychronization construct + code cleanup..
[OPNEMP, NVPTX] Fixed sychronization construct + code cleanup.
Mon, Jul 23, 6:56 AM
ABataev committed rL337691: [OPNEMP, NVPTX] Fixed sychronization construct + code cleanup..
[OPNEMP, NVPTX] Fixed sychronization construct + code cleanup.
Mon, Jul 23, 6:52 AM
ABataev closed D49564: [OPNEMP, NVPTX] Fixed sychronization construct + code cleanup..
Mon, Jul 23, 6:52 AM
ABataev added inline comments to D49564: [OPNEMP, NVPTX] Fixed sychronization construct + code cleanup..
Mon, Jul 23, 6:35 AM
ABataev added a comment to D44186: [OpenMP][libomptarget] New map interface: remove translation code and ensure proper alignment of struct members.

Yes, it was committed

Mon, Jul 23, 2:35 AM · Restricted Project

Jul 20 2018

ABataev updated the diff for D49564: [OPNEMP, NVPTX] Fixed sychronization construct + code cleanup..

Fixed SPMD mode barrier for initialized runtime.

Jul 20 2018, 1:26 PM
ABataev updated the diff for D49148: [DEBUGINFO] Disable unsupported debug info options for NVPTX target..

Added checks for all possible debug options.

Jul 20 2018, 1:04 PM
ABataev added inline comments to D49491: [RFC][VPlan, SLP] Add simple SLP analysis on top of VPlan..
Jul 20 2018, 12:58 PM
ABataev added inline comments to D49491: [RFC][VPlan, SLP] Add simple SLP analysis on top of VPlan..
Jul 20 2018, 10:53 AM
ABataev added inline comments to D49491: [RFC][VPlan, SLP] Add simple SLP analysis on top of VPlan..
Jul 20 2018, 7:20 AM

Jul 19 2018

ABataev created D49564: [OPNEMP, NVPTX] Fixed sychronization construct + code cleanup..
Jul 19 2018, 1:42 PM
ABataev committed rL337468: The patch adds support for the new map interface between clang and libomptarget..
The patch adds support for the new map interface between clang and libomptarget.
Jul 19 2018, 9:39 AM
ABataev committed rC337468: The patch adds support for the new map interface between clang and libomptarget..
The patch adds support for the new map interface between clang and libomptarget.
Jul 19 2018, 9:39 AM
ABataev 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:39 AM · Restricted Project
ABataev added a comment to D49510: [OpenMP][Clang] Usability improvements for OpenMP offloading.
  1. The new tool definitely requires an RFC at first.
  2. The patch is too big and should be split into several small patches
Jul 19 2018, 6:10 AM