In D60455#1467279, @aaron.ballman wrote:

In D60455#1464324, @Fznamznon wrote:

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?

In D60455#1468386, @Fznamznon wrote:

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?

LGTM. One minor comment.

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

LGTM. Thanks!

OpenCL C++ part looks good. Thanks!

LGTM. Thanks!

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.

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

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

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

Applied comments from @ABataev.

In D57768#1386941, @ABataev wrote:

In D57768#1386933, @bader wrote:

In D57768#1386924, @ABataev wrote:

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.

In D57768#1386924, @ABataev wrote:

This definitely requires a test.

In D57768#1386754, @Naghasan wrote:

LGTM

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

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

LGTM. Thanks!

Hi Sam,

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

Looks good to me. Thanks for working on this.

LGTM.

LGTM. Thanks!

In D39129#923235, @Anastasia wrote:

In D39129#902848, @bader wrote:

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

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

@arichardson, great work! Thanks a lot!

Applied comments.

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.

LGTM.

LGTM. Thanks!

In D36410#856716, @yaxunl wrote:

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

In D36410#855282, @Anastasia wrote:

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

LGTM, except one nit.

In D36327#840616, @yaxunl wrote:

In D36327#839809, @rjmccall wrote:

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.

LGTM! Thanks.

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

Hi Sam,

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

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.

LGTM. Thanks.

LGTM, thanks.

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.

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.

LGTM.
Just one minor comment.
Thanks.

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

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.

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.

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.

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.

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?

Please, split this patch into two parts:

1. Improve diagnostics on extension enabling.
2. Add missing sub_group_* built-in functions.