Page MenuHomePhabricator

bader (Alexey Bader)
User

Projects

User does not belong to any projects.

User Details

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

Recent Activity

Tue, Apr 16

bader added a comment to D60455: [SYCL] Add support for SYCL device attributes.

Applied comments from @aaron.ballman and @keryell

  • Introduced a C++11 and C2x style spelling in the clang namespace. I didn't find path to add two namespaces to attribute (like [[clang::sycl::device]]) so [[clang::sycl_device]] spelling is added.
  • Since both attributes have spellings and possible can be used as some "standard" outlining in Clang/LLVM I added documetation.
  • Added more test cases.
  • As @bader mentioned sycl_device can be used to mark functions, which are called from the different translation units so I added simple handling of this attribute in Sema.

I'm confused -- I thought @bader also said "...SYCL is not supposed to expose any non-standard extensions to a user." -- these attributes are not standards based (WG21 and WG14 have nothing to say about them), so are these attributes considered "non-standard extensions" or not?

Tue, Apr 16, 9:01 AM · Restricted Project
bader added a comment to D60455: [SYCL] Add support for SYCL device attributes.

Ok, my question is whether you are planning to duplicate the same logic as for OpenCL kernel which doesn't really seem like an ideal design choice. Is this the only difference then we can simply add an extra check for SYCL compilation mode in this template handling case. The overall interaction between OpenCL and SYCL implementation is still a very big unknown to me so it's not very easy to judge about the implementations details...

Of course, if nothing prevents us to re-use OpenCL kernel attribute for SYCL I assume it would be good idea.
But I'm thinking about the situation with https://reviews.llvm.org/D60454 . If we re-use OpenCL kernel attributes - we affected by OpenCL-related changes and OpenCL-related changes shouldn't violate SYCL semantics. Will it be usable for SYCL/OpenCL clang developers? @bader , what do you think about it?

Tue, Apr 16, 8:38 AM · Restricted Project

Mon, Apr 15

bader accepted D60455: [SYCL] Add support for SYCL device attributes.
Mon, Apr 15, 4:42 AM · Restricted Project
bader accepted D60455: [SYCL] Add support for SYCL device attributes.

LGTM. One minor comment.

Mon, Apr 15, 4:34 AM · Restricted Project
bader added inline comments to D60455: [SYCL] Add support for SYCL device attributes.
Mon, Apr 15, 2:47 AM · Restricted Project

Fri, Apr 12

bader accepted D60455: [SYCL] Add support for SYCL device attributes.
Fri, Apr 12, 7:43 AM · Restricted Project

Tue, Apr 9

bader added a comment to D60455: [SYCL] Add support for SYCL device attributes.

To give more context to the question, I'd like to clarify the use case of new attributes.

Tue, Apr 9, 1:51 PM · Restricted Project

Mar 25 2019

bader added inline comments to D59646: [PR40778][PR41157][OpenCL] Prevent implicit initialization of local address space objects.
Mar 25 2019, 5:50 AM · Restricted Project

Mar 24 2019

bader added inline comments to D59646: [PR40778][PR41157][OpenCL] Prevent implicit initialization of local address space objects.
Mar 24 2019, 1:40 PM · Restricted Project

Mar 19 2019

bader accepted D59544: [OpenCL] Minor improvements in default header testing.

LGTM. Thanks!

Mar 19 2019, 7:10 AM · Restricted Project

Mar 18 2019

bader accepted D59486: [OpenCL] Improve testing of default header in C++ mode.

OpenCL C++ part looks good. Thanks!

Mar 18 2019, 11:18 AM · Restricted Project
bader accepted D59492: [OpenCL] Allow variadic macros as Clang feature.

LGTM. Thanks!

Mar 18 2019, 8:18 AM
bader added a comment to D59486: [OpenCL] Improve testing of default header in C++ mode.

I see seven OpenCL C tests using -finclude-default-header option and AFAIK, only test/Driver/include-default-header.cl doesn't parse it. Could you also update OpenCL C tests?
I think clang/test/Headers/opencl-c-header.cl and test/Driver/include-default-header.cl fully cover this feature.
All other tests seems to use this option to simplify the test, but with additional cost on parsing this header.

Mar 18 2019, 7:05 AM · Restricted Project

Mar 15 2019

bader added a comment to D59219: [PR41007][OpenCL] Allow printf and toolchain reserved variadic functions in C++.

May be test/Driver/include-default-header.cl would make more sense to show the intent?

Mar 15 2019, 7:58 AM · Restricted Project

Mar 12 2019

bader added inline comments to D59219: [PR41007][OpenCL] Allow printf and toolchain reserved variadic functions in C++.
Mar 12 2019, 9:58 AM · Restricted Project

Feb 25 2019

bader closed D57768: [SYCL] Add clang front-end option to enable SYCL device compilation flow..

Committed: https://reviews.llvm.org/rL354773

Feb 25 2019, 3:53 AM · Restricted Project
bader committed rG3f62fa69a720: [SYCL] Add clang front-end option to enable SYCL device compilation flow. (authored by bader).
[SYCL] Add clang front-end option to enable SYCL device compilation flow.
Feb 25 2019, 3:50 AM
bader committed rC354773: [SYCL] Add clang front-end option to enable SYCL device compilation flow..
[SYCL] Add clang front-end option to enable SYCL device compilation flow.
Feb 25 2019, 3:49 AM
bader committed rL354773: [SYCL] Add clang front-end option to enable SYCL device compilation flow..
[SYCL] Add clang front-end option to enable SYCL device compilation flow.
Feb 25 2019, 3:49 AM

Feb 21 2019

bader updated the diff for D57768: [SYCL] Add clang front-end option to enable SYCL device compilation flow..

Add check that SYCL specific macro is not enabled by default.

Feb 21 2019, 6:42 AM · Restricted Project
bader retitled D57768: [SYCL] Add clang front-end option to enable SYCL device compilation flow. from [SYCL] Add SYCL device compilation flow. to [SYCL] Add clang front-end option to enable SYCL device compilation flow..
Feb 21 2019, 3:00 AM · Restricted Project
bader updated the diff for D57768: [SYCL] Add clang front-end option to enable SYCL device compilation flow..

Applied comments from @ABataev.

Feb 21 2019, 2:58 AM · Restricted Project

Feb 19 2019

bader committed rG24fa0c18e66c: [OpenCL] Change type of block pointer for OpenCL (authored by bader).
[OpenCL] Change type of block pointer for OpenCL
Feb 19 2019, 7:21 AM
bader committed rC354337: [OpenCL] Change type of block pointer for OpenCL.
[OpenCL] Change type of block pointer for OpenCL
Feb 19 2019, 7:21 AM
bader committed rL354337: [OpenCL] Change type of block pointer for OpenCL.
[OpenCL] Change type of block pointer for OpenCL
Feb 19 2019, 7:21 AM
bader closed D58277: [OpenCL] Change type of block pointer for OpenCL.
Feb 19 2019, 7:21 AM · Restricted Project, Restricted Project

Feb 11 2019

bader added a comment to D57768: [SYCL] Add clang front-end option to enable SYCL device compilation flow..

This definitely requires a test.

@ABataev, I tried to find some tests on similar -fcuda-is-device and -fopenmp-is-device options, but I wasn't able to find a dedicated test. Could you suggest some examples testing similar functionality, please?

There are several similar tests:
OpenMP/driver.c, Driver/openmp-offload.c, Driver/openmp-offload-gpu.c. There is no absolutely the same test for OpenMP, since OpenMP has mo similar req for the offloading.

Feb 11 2019, 7:08 AM · Restricted Project

Feb 6 2019

bader added a comment to D57768: [SYCL] Add clang front-end option to enable SYCL device compilation flow..

This definitely requires a test.

Feb 6 2019, 5:28 AM · Restricted Project
bader added a comment to D57768: [SYCL] Add clang front-end option to enable SYCL device compilation flow..

LGTM

Side note: might be good to also involve @Anastasia, as some of the future patches will overlap with OpenCL.

Feb 6 2019, 2:29 AM · Restricted Project

Feb 5 2019

bader created D57768: [SYCL] Add clang front-end option to enable SYCL device compilation flow..
Feb 5 2019, 10:28 AM · Restricted Project

Nov 28 2018

bader added inline comments to D54858: [OpenCL] Improve diagnostics for address spaces in template instantiation.
Nov 28 2018, 8:09 AM

Aug 24 2018

bader accepted D51212: [OpenCL][Docs] Release notes for OpenCL in Clang .
Aug 24 2018, 9:02 AM

Aug 17 2018

bader accepted D50259: [OpenCL] Disallow negative attribute arguments.

LGTM.
Please, remove "Change-Id: I910b5c077f5f29e02a1572d9202f0fdbea5280fd" from the log message - it not relevant to the project.

Aug 17 2018, 8:03 AM

Jun 20 2018

bader committed rL335103: [Sema] Allow creating types with multiple of the same addrspace..
[Sema] Allow creating types with multiple of the same addrspace.
Jun 20 2018, 1:36 AM
bader committed rC335103: [Sema] Allow creating types with multiple of the same addrspace..
[Sema] Allow creating types with multiple of the same addrspace.
Jun 20 2018, 1:36 AM
bader closed D47630: [Sema] Allow creating types with multiple of the same addrspace..
Jun 20 2018, 1:35 AM

Jun 7 2018

bader added inline comments to D47630: [Sema] Allow creating types with multiple of the same addrspace..
Jun 7 2018, 9:18 AM

Jun 1 2018

bader added inline comments to D47630: [Sema] Allow creating types with multiple of the same addrspace..
Jun 1 2018, 7:24 AM

May 23 2018

bader added inline comments to D46501: [OpenCL] Guard all half float usage based on cl_khr_fp16 extension.
May 23 2018, 4:18 AM · Restricted Project

Apr 27 2018

bader added inline comments to D46049: [OpenCL] Add constant address space to __func__ in AST.
Apr 27 2018, 6:14 AM

Apr 26 2018

bader added inline comments to D46049: [OpenCL] Add constant address space to __func__ in AST.
Apr 26 2018, 1:43 PM
bader accepted D46049: [OpenCL] Add constant address space to __func__ in AST.
Apr 26 2018, 10:38 AM

Apr 25 2018

bader added inline comments to D46049: [OpenCL] Add constant address space to __func__ in AST.
Apr 25 2018, 7:05 AM

Apr 11 2018

bader accepted D45363: [OpenCL] Added -std/-cl-std=CL2.2/CLC++.

LGTM. Thanks!

Apr 11 2018, 5:45 AM

Apr 6 2018

bader added inline comments to D45363: [OpenCL] Added -std/-cl-std=CL2.2/CLC++.
Apr 6 2018, 8:29 AM

Mar 6 2018

bader accepted D43783: [OpenCL] Remove block invoke function from emitted block literal struct.

Hi Sam,

Mar 6 2018, 2:39 AM

Feb 2 2018

bader added a reviewer for D42842: [MCJIT] Eliminate overhead in case of multiple subsequent calls to llvm::MCJIT::getFunctionAddress: lhames.
Feb 2 2018, 6:44 AM

Jan 26 2018

bader committed rL323522: [OpenCL] Add "cles_khr_int64" extension..
[OpenCL] Add "cles_khr_int64" extension.
Jan 26 2018, 3:51 AM
bader committed rC323522: [OpenCL] Add "cles_khr_int64" extension..
[OpenCL] Add "cles_khr_int64" extension.
Jan 26 2018, 3:51 AM
bader closed D42532: [OpenCL] Add "cles_khr_int64" extension..
Jan 26 2018, 3:51 AM

Jan 25 2018

bader accepted D42532: [OpenCL] Add "cles_khr_int64" extension..

LGTM, but I'd like Anastasia to approve.
Thanks.

Jan 25 2018, 6:25 AM
bader edited reviewers for D42532: [OpenCL] Add "cles_khr_int64" extension., added: Anastasia; removed: bader.
Jan 25 2018, 6:05 AM

Jan 21 2018

bader accepted D42307: [OpenCL][6.0.0 Release] Release notes for OpenCL in Clang.

Looks good to me. Thanks for working on this.

Jan 21 2018, 12:38 PM

Dec 18 2017

bader accepted D41344: Use unicode version of CreateFileMapping function..

LGTM.

Dec 18 2017, 7:34 AM
bader updated subscribers of D41344: Use unicode version of CreateFileMapping function..
Dec 18 2017, 7:30 AM

Nov 23 2017

bader accepted D39936: [OpenCL] Add extensions cl_intel_subgroups and cl_intel_subgroups_short.

LGTM. Thanks!

Nov 23 2017, 3:22 AM

Nov 15 2017

bader committed rL318290: [OpenCL] Fix code generation of function-scope constant samplers..
[OpenCL] Fix code generation of function-scope constant samplers.
Nov 15 2017, 3:38 AM
bader closed D34342: [OpenCL] Fix code generation of function-scope constant samplers. by committing rL318290: [OpenCL] Fix code generation of function-scope constant samplers..
Nov 15 2017, 3:38 AM

Nov 14 2017

bader added a comment to D39129: [OpenCL] Fix generation of constant address space sampler in function scope.

@Anastasia, during the discussion of similar fix (https://reviews.llvm.org/D34342).

I found another bug in the CodeGen library. Do you have time to fix it too?
Here is the reproducer from the old review request:

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

I am getting an error currently: error: internal error: could not emit constant value "abstractly". It seems very weird and also doesn't seem to be related to constant AS change. I think this should be allowed to compile though? I can try to investigate but since I don't know the time frame yet and it doesn't seem to be related to constant AS I would prefer to do it as a separate change.

Nov 14 2017, 3:16 AM

Oct 21 2017

bader added a comment to D39129: [OpenCL] Fix generation of constant address space sampler in function scope.

@Anastasia, during the discussion of similar fix (https://reviews.llvm.org/D34342).

Oct 21 2017, 6:03 AM

Oct 12 2017

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

@arichardson, great work! Thanks a lot!

Oct 12 2017, 10:20 AM

Oct 11 2017

bader committed rL315453: [OpenCL] Allow function declaration with empty argument list..
[OpenCL] Allow function declaration with empty argument list.
Oct 11 2017, 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..
Oct 11 2017, 4:16 AM

Oct 10 2017

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

Applied comments.

Oct 10 2017, 11:31 AM

Oct 9 2017

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.

Oct 9 2017, 8:37 AM

Oct 3 2017

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

LGTM.

Oct 3 2017, 5:21 AM

Sep 25 2017

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

LGTM. Thanks!

Sep 25 2017, 7:23 AM

Sep 22 2017

bader added inline comments to D37804: [OpenCL] Handle address space conversion while setting type alignment.
Sep 22 2017, 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