bader (Alexey Bader)
User

Projects

User does not belong to any projects.

User Details

User Since
Apr 23 2014, 2:13 AM (182 w, 1 d)

Recent Activity

Thu, Oct 12

bader accepted D38816: Convert clang::LangAS to a strongly typed enum.

@arichardson, great work! Thanks a lot!

Thu, Oct 12, 10:20 AM

Wed, Oct 11

bader committed rL315453: [OpenCL] Allow function declaration with empty argument list..
[OpenCL] Allow function declaration with empty argument list.
Wed, Oct 11, 4:17 AM
bader closed D33681: [OpenCL] Allow function declaration with empty argument list. by committing rL315453: [OpenCL] Allow function declaration with empty argument list..
Wed, Oct 11, 4:16 AM

Tue, Oct 10

bader updated the diff for D33681: [OpenCL] Allow function declaration with empty argument list..

Applied comments.

Tue, Oct 10, 11:31 AM

Mon, Oct 9

bader updated the diff for D33681: [OpenCL] Allow function declaration with empty argument list..

Fix parsing of function declarations with empty parameter list initializers.
This change should fix the failure caught by the buildbot.
Sorry for the long delay.

Mon, Oct 9, 8:37 AM

Tue, Oct 3

bader accepted D38463: [OpenCL] Fix checking of vector type casting.

LGTM.

Tue, Oct 3, 5:21 AM

Mon, Sep 25

bader accepted D37804: [OpenCL] Handle address space conversion while setting type alignment.

LGTM. Thanks!

Mon, Sep 25, 7:23 AM

Fri, Sep 22

bader added inline comments to D37804: [OpenCL] Handle address space conversion while setting type alignment.
Fri, Sep 22, 11:01 AM

Aug 30 2017

bader added a comment to D36410: [OpenCL] Handle taking address of block captures.

The captured variable is still passed by value. The address taking is on the duplicate of the captured variable, not on the original variable.

Aug 30 2017, 11:50 AM

Aug 29 2017

bader added a comment to D36410: [OpenCL] Handle taking address of block captures.

Ok, I will update it to be implicitly generic then. Just to be sure, @bader do you agree on this too?

Aug 29 2017, 5:43 AM

Aug 21 2017

bader accepted D36951: [OpenCL][5.0.0 Release] Release notes for OpenCL in Clang.

LGTM, except one nit.

Aug 21 2017, 4:55 AM

Aug 14 2017

bader added a comment to D36327: [OpenCL] Allow targets emit optimized pipe functions for power of 2 type sizes.

Could you just implement this in SimplifyLibCalls? I assume there's some way to fill in TargetLibraryInfo appropriately for a platform. Is that too late for your linking requirements?

Both the optimized and generic versions of __read_pipe function contains call of other library functions and are complicate enough not to be generated programmatically. amdgpu target does not have the capability to link in library code after LLVM codegen. The linking has to be done before SimplifyLibCalls.

Aug 14 2017, 7:43 AM
bader accepted D36676: Remove -finclude-default-header in OpenCL atomic tests.

LGTM! Thanks.

Aug 14 2017, 7:06 AM

Aug 10 2017

bader added inline comments to D32896: [OpenCL] Make CLK_NULL_RESERVE_ID invalid reserve id..
Aug 10 2017, 7:52 AM

Aug 8 2017

bader added a comment to D36327: [OpenCL] Allow targets emit optimized pipe functions for power of 2 type sizes.

@rsmith do you have an opinion on what would be the right place for the kind of proposed optimization?
It looks like it can be implemented as target independent optimization, acting only for target with specified properties - in this case target must provide required built-in functions.

Aug 8 2017, 11:46 AM
bader added a reviewer for D36327: [OpenCL] Allow targets emit optimized pipe functions for power of 2 type sizes: rsmith.
Aug 8 2017, 11:35 AM

Aug 7 2017

bader added a comment to D36327: [OpenCL] Allow targets emit optimized pipe functions for power of 2 type sizes.

Hi Sam,

Aug 7 2017, 3:27 AM

Aug 4 2017

bader accepted D28691: Add OpenCL 2.0 atomic builtin functions as Clang builtin.
Aug 4 2017, 8:11 AM

Aug 2 2017

bader accepted D36259: [OpenCL] Remove extra select functions from opencl-c.h.

Out of curiosity: how did you detect this?
Can we use the same approach for writing tests?

Aug 2 2017, 11:47 PM

Jul 31 2017

bader added inline comments to D36044: [OpenCL] -cl-ext option can overwrite OpenCL features imported from a module.
Jul 31 2017, 7:41 AM

Jul 27 2017

bader accepted D33945: [OpenCL] Add support for missing sub_group functions..

Thanks!
Overall the patch looks good, but I would suggest splitting it into three commits (as they seems to be independent):

  1. [OpenCL] Check that cl_khr_subgroups pragma is enabled if respective extension is used.
  2. [OpenCL] Add support for missing sub_group functions.
  3. [OpenCL] Fix return type for reserve pipe built-ins.
Jul 27 2017, 7:07 AM

Jul 14 2017

bader accepted D35420: [OpenCL] Fix access qualifiers metadata for kernel arguments with typedef.

LGTM. Thanks.

Jul 14 2017, 11:41 AM

Jul 11 2017

bader accepted D35082: [OpenCL] Add LangAS::opencl_private to represent private address space in AST.

LGTM, thanks.

Jul 11 2017, 3:41 AM

Jul 4 2017

bader added a comment to D34948: [OpenCL] Generalise err_opencl_enqueue_kernel_expected_type to be used with other builtins.

Do you have another built-in in mind which can use this diagnostic message?
If so, it would make sense to re-use it in the same patch.

This is split off from https://reviews.llvm.org/D33945, which I will be rebasing/re-uploading once this patch is committed.

Jul 4 2017, 3:06 AM

Jul 3 2017

bader accepted D34948: [OpenCL] Generalise err_opencl_enqueue_kernel_expected_type to be used with other builtins.

LGTM.
Do you have another built-in in mind which can use this diagnostic message?
If so, it would make sense to re-use it in the same patch.

Jul 3 2017, 9:00 AM

Jun 30 2017

bader accepted D34871: [OpenCL] Add function name to extension diagnostic.

LGTM.
Just one minor comment.
Thanks.

Jun 30 2017, 2:48 AM

Jun 29 2017

bader added a comment to D33681: [OpenCL] Allow function declaration with empty argument list..

@chapuni, thanks for taking care of this. I'll take a look.

Jun 29 2017, 5:05 AM
bader committed rL306653: [OpenCL] Allow function declaration with empty argument list..
[OpenCL] Allow function declaration with empty argument list.
Jun 29 2017, 1:44 AM
bader closed D33681: [OpenCL] Allow function declaration with empty argument list..
Jun 29 2017, 1:44 AM

Jun 27 2017

bader added a comment to D34342: [OpenCL] Fix code generation of function-scope constant samplers..

Note: get_sampler_initializer from my test case returns integer, not a sampler, but having function is not relevant to the problem.
Here is a bit simplified test case without function calls that still reproduces the problem:

kernel void foo(int sampler_init_value) {
  const sampler_t const_smp_func_init = sampler_init_value;
}

The problem is in the function that handles sampler initialization with integer.
There should no problems with the case you provided as it doesn't require additional function call injection.

Jun 27 2017, 10:03 AM

Jun 23 2017

bader added a comment to D33681: [OpenCL] Allow function declaration with empty argument list..

Although this patch is already accepted, I'd like to confirm that I can commit the latest update with changes in test/SemaOpenCL/invalid-pipes-cl2.0.cl.

Jun 23 2017, 7:43 AM

Jun 21 2017

bader updated the diff for D34342: [OpenCL] Fix code generation of function-scope constant samplers..

Added test case reproducing the issue described in the description.
Removed test cases from test/SemaOpenCL/sampler_t.cl covered by test/CodeGenOpenCL/sampler.cl.

Jun 21 2017, 10:39 AM
bader added a comment to D34342: [OpenCL] Fix code generation of function-scope constant samplers..

I think I found the way to reproduce the issue. There are two pre-requisites: sampler must be declared in the constant address space and clang must be build in Debug mode.
I'll fix the test and add the test cases from test/SemaOpenCL/sampler_t.cl as mentioned earlier.

Jun 21 2017, 8:54 AM

Jun 20 2017

bader abandoned D34342: [OpenCL] Fix code generation of function-scope constant samplers..
Jun 20 2017, 11:43 PM
bader added a comment to D34342: [OpenCL] Fix code generation of function-scope constant samplers..

Wow...
Nice catch.
For some reason I can't reproduce the problem neither.
I will drop source code change, but I'd like to modify the test anyway.
https://reviews.llvm.org/rL277024 added test/CodeGenOpenCL/sampler.cl and modified test/SemaOpenCL/sampler_t.cl.
I want to move the parts of the test/SemaOpenCL/sampler_t.cl, which are supposed to pass semantic analysis to test/CodeGenOpenCL/sampler.cl and validate the LLVM IR after code generation.
Are you OK with this?

Jun 20 2017, 9:05 AM
bader committed rL305796: [OpenCL] Fix OpenCL and SPIR version metadata generation..
[OpenCL] Fix OpenCL and SPIR version metadata generation.
Jun 20 2017, 7:31 AM
bader closed D34235: [OpenCL] Fix OpenCL and SPIR version metadata generation..
Jun 20 2017, 7:30 AM

Jun 19 2017

bader added inline comments to D34235: [OpenCL] Fix OpenCL and SPIR version metadata generation..
Jun 19 2017, 10:16 AM
bader added inline comments to D34342: [OpenCL] Fix code generation of function-scope constant samplers..
Jun 19 2017, 9:01 AM
bader requested changes to D33945: [OpenCL] Add support for missing sub_group functions..

Please, split this patch into two parts:

  1. Improve diagnostics on extension enabling.
  2. Add missing sub_group_* built-in functions.
Jun 19 2017, 4:09 AM
bader created D34342: [OpenCL] Fix code generation of function-scope constant samplers..
Jun 19 2017, 1:41 AM

Jun 16 2017

bader added inline comments to D33989: [OpenCL] Allow targets to select address space per type.
Jun 16 2017, 5:02 AM
bader updated the diff for D34235: [OpenCL] Fix OpenCL and SPIR version metadata generation..

Applied Akira's comment: re-used cached Int32Ty type.

Jun 16 2017, 12:51 AM

Jun 15 2017

bader created D34235: [OpenCL] Fix OpenCL and SPIR version metadata generation..
Jun 15 2017, 4:46 AM

Jun 14 2017

bader added inline comments to D33989: [OpenCL] Allow targets to select address space per type.
Jun 14 2017, 2:22 AM

Jun 13 2017

bader retitled D33681: [OpenCL] Allow function declaration with empty argument list. from Allow function declaration with empty argument list. to [OpenCL] Allow function declaration with empty argument list..
Jun 13 2017, 5:47 AM
bader updated the diff for D33681: [OpenCL] Allow function declaration with empty argument list..

Update one more test missed in the first version of the patch.

Jun 13 2017, 5:45 AM

Jun 2 2017

bader closed D33821: [OpenCL] Harden function pointer diagnostics. by committing rL304575: [OpenCL] Harden function pointer diagnostics..
Jun 2 2017, 11:09 AM
bader committed rL304575: [OpenCL] Harden function pointer diagnostics..
[OpenCL] Harden function pointer diagnostics.
Jun 2 2017, 11:09 AM
bader updated subscribers of D33681: [OpenCL] Allow function declaration with empty argument list..
Jun 2 2017, 10:54 AM
bader created D33821: [OpenCL] Harden function pointer diagnostics..
Jun 2 2017, 3:31 AM
bader added a comment to D33681: [OpenCL] Allow function declaration with empty argument list..

Sorry for the delay.
It will also align the behavior with function definition.
For instance the following function definition is allowed.

__local int *allocateLocalMemory() { return 0; }
Jun 2 2017, 1:31 AM

May 30 2017

bader created D33681: [OpenCL] Allow function declaration with empty argument list..
May 30 2017, 8:49 AM

May 26 2017

bader added a comment to D33353: [OpenCL] An error shall occur if any scalar operand has greater rank than the type of the vector element.

This issue probably can be fixed by setting SPIR target architecture or just enabling cl_khr_fp64 extension.
Unfortunately, I can't check it today.
Could you revert Egor's commit, please?

May 26 2017, 8:43 AM

May 23 2017

bader added inline comments to D32896: [OpenCL] Make CLK_NULL_RESERVE_ID invalid reserve id..
May 23 2017, 3:53 AM

May 22 2017

bader added inline comments to D32896: [OpenCL] Make CLK_NULL_RESERVE_ID invalid reserve id..
May 22 2017, 7:58 AM
bader added inline comments to D32896: [OpenCL] Make CLK_NULL_RESERVE_ID invalid reserve id..
May 22 2017, 3:42 AM

May 17 2017

bader added inline comments to D32896: [OpenCL] Make CLK_NULL_RESERVE_ID invalid reserve id..
May 17 2017, 8:40 AM
bader edited reviewers for D32896: [OpenCL] Make CLK_NULL_RESERVE_ID invalid reserve id., added: yaxunl; removed: cfe-commits.
May 17 2017, 8:40 AM

Mar 28 2017

bader accepted D31397: [Bug 25404] Fix crash on typedef in OpenCL 2.0.

I think my topics are not related to the bug 25404 and can be handled separately.
Thanks!

Mar 28 2017, 8:35 AM
bader added inline comments to D31397: [Bug 25404] Fix crash on typedef in OpenCL 2.0.
Mar 28 2017, 5:20 AM
bader added a comment to D31397: [Bug 25404] Fix crash on typedef in OpenCL 2.0.

Anastasia, could you generate patch with full context, please?

Mar 28 2017, 3:08 AM

Mar 16 2017

bader removed a reviewer for D30816: [OpenCL] Added implicit conversion rank for overloading functions with vector data type in OpenCL: bader.
Mar 16 2017, 9:01 AM
bader removed a reviewer for D30937: [OpenCL] Added diagnostic for checking length of vector: bader.
Mar 16 2017, 8:53 AM

Mar 15 2017

bader added inline comments to D28136: [OpenCL] Implement as_type operator as alias of __builtin_astype..
Mar 15 2017, 9:26 AM
bader added a comment to D28136: [OpenCL] Implement as_type operator as alias of __builtin_astype..

From all the above arguments, I feel like the right approach would be to implement it as Clang builtin which closely matches the operator semantic in my opinion. We could of course reuse the implementation of __bultin_astype to avoid unnecessary extra work and code duplication.

Using the macro seems to me more like a workaround solution and overloaded functions don't seem to be entirely the right thing either. What do you think about it?

Mar 15 2017, 3:35 AM
bader added inline comments to D30805: [OpenCL] Add intel_reqd_sub_group_size attribute support.
Mar 15 2017, 1:58 AM

Mar 14 2017

bader added inline comments to D30643: [OpenCL] Extended diagnostics for atomic initialization.
Mar 14 2017, 5:51 AM
bader added inline comments to D30816: [OpenCL] Added implicit conversion rank for overloading functions with vector data type in OpenCL.
Mar 14 2017, 5:38 AM

Mar 13 2017

bader added inline comments to D30643: [OpenCL] Extended diagnostics for atomic initialization.
Mar 13 2017, 5:32 AM
bader added inline comments to D30643: [OpenCL] Extended diagnostics for atomic initialization.
Mar 13 2017, 3:26 AM

Mar 10 2017

bader added a comment to D28136: [OpenCL] Implement as_type operator as alias of __builtin_astype..

Why do you think this is a bug? It seems to follow standard behavior in C to promote char to int if required. Just like if you would have a C code:

int as_int(int i);
void foo() {
    char src = 1;
    int dst = as_int(src);
}

This code would complie and the same exactly IR would be generated.

Mar 10 2017, 8:56 AM
bader accepted D30805: [OpenCL] Add intel_reqd_sub_group_size attribute support.

LGTM.
Just a few minor issues.

Mar 10 2017, 3:00 AM

Jan 29 2017

bader committed rL293467: [LanRef] Fix typo in getelementptr example..
[LanRef] Fix typo in getelementptr example.
Jan 29 2017, 11:50 PM
bader closed D29009: [LanRef] Fix typo in getelementptr example..
Jan 29 2017, 11:50 PM

Jan 23 2017

bader created D29009: [LanRef] Fix typo in getelementptr example..
Jan 23 2017, 3:29 AM

Jan 18 2017

bader added inline comments to D28080: [Docs][OpenCL] Added OpenCL feature description to user manual..
Jan 18 2017, 11:49 AM

Jan 11 2017

bader accepted D28080: [Docs][OpenCL] Added OpenCL feature description to user manual..

LGTM.
Thanks!

Jan 11 2017, 6:15 AM
bader added a comment to D28080: [Docs][OpenCL] Added OpenCL feature description to user manual..

Thanks for working on this!
Looks good, except a few pedantic notes.

Jan 11 2017, 1:09 AM

Dec 7 2016

bader added a comment to D27403: [OpenCL] Added a LIT test for ensuring address space mangling is done the same both in OpenCL1.2 and OpenCL2.0..

@echuraev, please, request commit access as described here: http://llvm.org/docs/DeveloperPolicy.html#obtaining-commit-access

Dec 7 2016, 12:56 AM
bader committed rL288891: [OpenCL] Added a LIT test for ensuring address space mangling is done the same….
[OpenCL] Added a LIT test for ensuring address space mangling is done the same…
Dec 7 2016, 12:54 AM
bader closed D27403: [OpenCL] Added a LIT test for ensuring address space mangling is done the same both in OpenCL1.2 and OpenCL2.0. by committing rL288891: [OpenCL] Added a LIT test for ensuring address space mangling is done the same….
Dec 7 2016, 12:53 AM
bader committed rL288890: [OpenCL] Fix SPIR version generation..
[OpenCL] Fix SPIR version generation.
Dec 7 2016, 12:48 AM
bader closed D27300: [OpenCL] Fix SPIR version generation. by committing rL288890: [OpenCL] Fix SPIR version generation..
Dec 7 2016, 12:48 AM

Dec 6 2016

bader added a comment to D27334: [OpenCL] Ambiguous function call..

This change seems to modify normal C behavior again. Is there any strong motivation for doing this and if yes could it be done generically with C?

Motivation:

// Non-portable OpenCL 1.2 code 
__kernel void foo(global float* out) {
  out[get_global_id(0)] = sin(get_global_id(0));
}

This program compiles fine on OpenCL platform w/o doubles support and fails otherwise.
If OpenCL driver supports doubles it provides at least two versions of 'sin' built-in math function and compiler will not be able to choose the right one for 'size_t' argument.
The goal of this warning is to let OpenCL developer know about potential issues with code portability.

I would argue this improves the portability much as it can also be misleading in some situations (because it refers to a potentially hypothetical problem). For example there can be builtin functions that only have a float parameter (without a double version of it). This is for example the case with read_image functions that take a float coordinate value between 0 and 1. Unfortunately this warning won't be triggered on read_image functions because there is an overload candidate with an int type of the same parameter too. But we can't exclude this situations to appear in the future or from some vendor extensions or even custom OpenCL code.

Dec 6 2016, 3:34 AM

Dec 4 2016

bader added a comment to D27334: [OpenCL] Ambiguous function call..

This change seems to modify normal C behavior again. Is there any strong motivation for doing this and if yes could it be done generically with C?

Dec 4 2016, 12:51 AM

Nov 29 2016

bader committed rL288126: [OpenCL] Prohibit using reserve_id_t in program scope..
[OpenCL] Prohibit using reserve_id_t in program scope.
Nov 29 2016, 2:31 AM
bader closed D27099: [OpenCL] Prohibit using reserve_id_t in program scope. by committing rL288126: [OpenCL] Prohibit using reserve_id_t in program scope..
Nov 29 2016, 2:31 AM

Nov 24 2016

bader removed a reviewer for D27049: [OpenCL] Refactor out ReadPipe/WritePipe: bader.
Nov 24 2016, 1:52 AM

Nov 17 2016

bader updated subscribers of D26746: [OpenCL] Split PipeType into ReadPipe/WritePipe.
Nov 17 2016, 3:35 AM
bader accepted D26746: [OpenCL] Split PipeType into ReadPipe/WritePipe.

LGTM. Thanks!
A few minor comments regarding outdated comments and style.

Nov 17 2016, 3:34 AM

Nov 16 2016

bader removed a reviewer for D26735: [OpenCL] Disable && (address of label) GNU extension for OpenCL: cfe-commits.
Nov 16 2016, 1:57 AM
bader added a reviewer for D26735: [OpenCL] Disable && (address of label) GNU extension for OpenCL: Anastasia.
Nov 16 2016, 1:56 AM

Nov 2 2016

bader added a comment to D26157: [OpenCL] always use SPIR address spaces for kernel_arg_addr_space MD.

@pekka.jaaskelainen, I have related question: what is the problem with retaining OpenCL address space information in LLVM IR?
My understanding is that target CodeGen can ignore this information for the machines with 'flat' address space model.
On the other hand I would expect this information be useful for the Optimizer to resolve pointer aliasing.
Does it make any sense to keep OpenCL address space information in LLVM IR generated for the machines with 'flat' address space (e.g. CPU)?

Nov 2 2016, 4:21 AM

Nov 1 2016

bader committed rL285700: [OpenCL] Override supported OpenCL extensions with -cl-ext option.
[OpenCL] Override supported OpenCL extensions with -cl-ext option
Nov 1 2016, 9:00 AM
bader closed D23712: [OpenCL] Override supported OpenCL extensions with -cl-ext option by committing rL285700: [OpenCL] Override supported OpenCL extensions with -cl-ext option.
Nov 1 2016, 9:00 AM

Oct 31 2016

bader added a comment to D25305: [OpenCL] Setting constant address space for array initializers.

Committed at 285557 with updated tests.

Oct 31 2016, 3:36 AM
bader committed rL285557: [OpenCL] Setting constant address space for array initializers.
[OpenCL] Setting constant address space for array initializers
Oct 31 2016, 3:36 AM

Oct 25 2016

bader added inline comments to D25305: [OpenCL] Setting constant address space for array initializers.
Oct 25 2016, 3:25 AM

Oct 24 2016

bader added inline comments to D25305: [OpenCL] Setting constant address space for array initializers.
Oct 24 2016, 3:37 AM