[SYCL] Add sycl_kernel attribute for accelerated code outlining

[bader]

Summary

SYCL is single source offload programming model relying on compiler to
separate device code (i.e. offloaded to an accelerator) from the code
executed on the host.

Here is code example of the SYCL program to demonstrate compiler
outlining work:

int foo(int x) { return ++x; }
int bar(int x) { throw std::exception("CPU code only!"); }
...
using namespace cl::sycl;
queue Q;
buffer<int, 1> a(range<1>{1024});
Q.submit([&](handler& cgh) {
  auto A = a.get_access<access::mode::write>(cgh);
  cgh.parallel_for<init_a>(range<1>{1024}, [=](id<1> index) {
    A[index] = index[0] * 2 + index[1] + foo(42);
  });
}
...

SYCL device compiler must compile lambda expression passed to
cl::sycl::handler::parallel_for method and function foo called from this
lambda expression for an "accelerator". SYCL device compiler also must
ignore bar function as it's not required for offloaded code execution.

This patch adds the sycl_kernel attribute, which is used to mark code
passed to cl::sycl::handler::parallel_for as "accelerated code".

Attribute must be applied to function templates which parameters include
at least "kernel name" and "kernel function object". These parameters
will be used to establish an ABI between the host application and
offloaded part.

[bader]

clang/lib/Sema/SemaSYCL.cpp

23 ↗

(On Diff #200658)

I think this is also preventing traditional linking of translation units.

Could you elaborate more on this topic, please?
What do you mean by "traditional linking of translation units" and what exactly "is preventing" it?
Do you compare with the linking of regular C++ code (i.e. which do not split into host and device code)?
If so, SYCL is different from this model and more similar to CUDA/OpenMP models, which also skip "linking" of irrelevant part (e.g. host code is not linked by the device compiler).
Mariya added Justin (@jlebar) and Alexey (@ABataev), who work on single-source programming models to make them aware and provide feedback if any.

[Anastasia]

In D60455#1518178, @bader wrote:

In D60455#1513499, @Anastasia wrote:

Design question: since you are not aware what functions are to be run on a device while parsing them, at what point do you plan to diagnose the invalid behavior from the standard C++ i.e. using function pointers in kernel code is likely to cause issues?

We are going to use DeviceDiagBuilder and related infrastructure implemented in Clang to diagnose device side code errors in OpenMP/CUDA modes.
More details are in the comments here:
https://clang.llvm.org/doxygen/classclang_1_1Sema_1_1DeviceDiagBuilder.html#details

Just a thought, if you parse host code first and provide the device outlining information to the device compilation phase would you then be able to reuse more parsing functionality from OpenCL?

Also do you need to outline the data structures too? For example classes used on device are not allowed to have virtual function.

Yes. This restriction is already implemented in our code base on GitHub.

Cool, is it implemented in SemaSYCL.cpp too?

clang/include/clang/Basic/Attr.td

1076

Sema part is mostly not relevant for SYCL mode because SYCL API do not allow the cases currently detected by clang (e.g. constant address space variable declaration in OpenCL kernel scope, naming OpenCL kernel main, etc).

Would you mind pointing me to your impl of those?

A couple of check that might be useful are:

void return type for kernel functions
kernel can't be static function

and some of the checks are harmful for proposed implementation (e.g. kernels can't be template functions).

@Anastasia, @keryell, @agozillon and @aaron.ballman need to agree if this sufficient to justify the re-use of OpenCL kernel attribute.
Let me know if you need any additional information to make a decision.

Ok, if from ~20 occurrences in the source code you will be able to reuse only just 2 it doesn't seem like it's worth to share __kernel attribute.

clang/lib/Sema/SemaSYCL.cpp

23 ↗

(On Diff #200658)

Yes indeed, I mean linking of modules in C/C++ even though it doesn't necessarily mean linking of object files. So you don't plan to support SYCL_EXTERNAL in clang?

In CUDA the functions executed on device are annotated manually using __device__ hence separate translation units can specify external device function... although I don't know if CUDA implementation in clang support this.

I guess OpenMP is allowed to fall back to run on host?

[Fznamznon]

clang/test/SemaSYCL/device-attributes-on-non-sycl.cpp

1 ↗

(On Diff #200658)

Sorry for confusion. The C++ features used in SYCL are a subset of the C++11 standard features.
I will add -std=c++11 key to run line to avoid such confusion in future.

[Fznamznon]

Applied comments from @Anastasia

• Added documentation for sycl_kernel function
• Added comments to Sema.h
• Added -std=c++11 to test run lines

[Anastasia]

clang/include/clang/Basic/Attr.td

1076

Undocumented -> SYCLKernelDocs

clang/include/clang/Basic/AttrDocs.td

269

The example doesn't demonstrate the use of the attribute.

It explains how it is used by the toolchain only!

May be @aaron.ballman can help here as I am not sure what the format should be.

clang/lib/Parse/ParseAST.cpp

171 ↗

(On Diff #201641)

Do you also need to prevent generation of non-device functions somehow?

clang/lib/Sema/SemaSYCL.cpp

23 ↗

(On Diff #200658)

Ping!

I would suggest to document it a bit more including any current limitations/assumption that you can mark under FIXME i.e. does your code handle lambdas yet, what if lambdas are used in function parameters, etc...

clang/lib/Sema/SemaTemplateInstantiateDecl.cpp

5520 ↗

(On Diff #201641)

May be this should go into a helper function as it seems to be now a bigger chunk of code that is repeated?

Although, I am not very familiar with this code. You can try to get someone to review who has contributed to this more recently.

clang/test/CodeGenSYCL/device-functions.cpp

24 ↗

(On Diff #201641)

I can't see where the SPIR calling convention is currently set for SYCL?

clang/test/SemaSYCL/device-attributes-on-non-sycl.cpp

3 ↗

(On Diff #201641)

I don't think this comment is necessary.

[Fznamznon]

clang/include/clang/Basic/Attr.td

1076

Oh, Thank you for that!

clang/lib/Parse/ParseAST.cpp

171 ↗

(On Diff #201641)

I think It's already prevented by change to CodeGenModule.cpp in this patch. CodeGen just ignores declarations without SYCL device attribute now.

clang/lib/Sema/SemaSYCL.cpp

23 ↗

(On Diff #200658)

Oh, sorry, I missed this comment when I updated patch last time.
Could you please advise in which form I can document it?

clang/lib/Sema/SemaTemplateInstantiateDecl.cpp

5520 ↗

(On Diff #201641)

I think this chunk of code seems big because of big repeated comment.

clang/test/CodeGenSYCL/device-functions.cpp

24 ↗

(On Diff #201641)

If I understand correct it's set automatically on AST level because we use SPIR-based triple for device code. Only in case of C++ methods clang doesn't set SPIR calling convention. We did a modification in our codebase to get SPIR calling convention for C++ methods too (available here )

[Anastasia]

clang/test/CodeGenSYCL/device-functions.cpp

24 ↗

(On Diff #201641)

Ok and what happens if some other target is used - not SPIR?

[Fznamznon]

clang/test/CodeGenSYCL/device-functions.cpp

24 ↗

(On Diff #201641)

There will be no SPIR calling convention for device functions.

[Anastasia]

clang/test/CodeGenSYCL/device-functions.cpp

24 ↗

(On Diff #201641)

Just FYI at some point we generalized SPIR calling convention to be used for kernels irrespective from target by default (see TargetCodeGenInfo::getOpenCLKernelCallingConv). Not sure if it might make sense to do for SYCL device functions too. I am not saying it belongs to this patch though.

[Fznamznon]

Applied comments from @Anastasia

• Added link to documentation for sycl_device attribute
• Removed redundant comment from test

@Anastasia, do you have additional comments?

[Fznamznon]

@aaron.ballman , please let me know if you have additional comments/suggestions. If not, could you please accept this revision?

[Fznamznon]

Ping.

[aaron.ballman]

Most of the comments are about minor nits like grammar and coding conventions, but I did have some questions regarding what kinds of functions the sycl_kernel attribute gets applied to. Also, I'd like to see some additional tests that demonstrate the sycl device attribute is being implicitly created on the proper declarations as expected (can probably do that using -ast-dump and checking to see if the right functions have the attribute attached).

clang/include/clang/Basic/AttrDocs.td

259

is SYCL "kernel function" -> is a SYCL "kernel function"

260

SYCL -> A SYCL

261

Kernel is a -> A kernel is a

263–264

This doesn't really demonstrate the need for the attribute -- the attribute is never shown in the code example. I'd prefer an example that shows when and how a user would write this attribute.

278

called SYLC -> called a SYLC

280

use sycl_kernel -> use the sycl_kernel

281

as SYCL -> as a SYCL
Compiler is supposed to -> The compiler will

284

In this code example compiler is supposed to add "foo" function -> In this code example, the compiler will add the "foo" function

clang/include/clang/Sema/Sema.h

11182 ↗

(On Diff #203785)

Function -> function

11183 ↗

(On Diff #203785)

In SYCL, when we generate device code, we don't know

11184 ↗

(On Diff #203785)

we emit sycl kernels, so we add device

11189 ↗

(On Diff #203785)

adds the function declaration

11190 ↗

(On Diff #203785)

Should be named addSyclDeviceFunc() per coding standards. Similar for the other new functions.

11194 ↗

(On Diff #203785)

Don't repeat the function name in the comments, please. Also, rather than returning a concrete SmallVector<>, I think it would be more natural to return a SmallVectorImpl so that callers don't have to contend with the explicit size. There should also be a const overload for this function.

11197 ↗

(On Diff #203785)

Constructs a SYCL kernel that is compatible with OpenCL from the SYCL "kernel

11200–11201 ↗

(On Diff #203785)

Marks all functions accessible from SYCL kernels with the SYCL device attribute

clang/lib/CodeGen/CodeGenModule.cpp

2410 ↗

(On Diff #203785)

with the SYCL device attribute

2412–2415 ↗

(On Diff #203785)

These if statements can be combined.

2533 ↗

(On Diff #203785)

Missing a full stop at the end of the comment.

clang/lib/Sema/SemaSYCL.cpp

14 ↗

(On Diff #203785)

This include doesn't seem to be necessary?

23 ↗

(On Diff #203785)

e does not use our usual naming conventions.

41 ↗

(On Diff #203785)

Spurious whitespace can be removed.

44 ↗

(On Diff #203785)

Elide braces.

48 ↗

(On Diff #203785)

Don't use auto as the type is not spelled out in the initialization.

52 ↗

(On Diff #203785)

Elide braces.

68 ↗

(On Diff #203785)

with the SYCL device attribute

70 ↗

(On Diff #203785)

elt -> Elt per naming conventions

71 ↗

(On Diff #203785)

auto * since the type is spelled out in the initialization.

73 ↗

(On Diff #203785)

Elide braces

clang/lib/Sema/SemaTemplateInstantiateDecl.cpp

5523 ↗

(On Diff #203785)

for the SYCL kernel attribute

5525 ↗

(On Diff #203785)

Elide braces

5537 ↗

(On Diff #203785)

for the SYCL kernel attribute

5539 ↗

(On Diff #203785)

Elide braces

clang/test/SemaSYCL/device-attributes-on-non-sycl.cpp

4 ↗

(On Diff #203785)

#ifndef ?

11 ↗

(On Diff #203785)

I'd prefer to spell this with __attribute__, same in the other test

clang/test/SemaSYCL/device-attributes.cpp

3 ↗

(On Diff #203785)

I'd like to see some more tests covering less obvious scenarios. Can I add this attribute to a lambda? What about a member function? How does it work with virtual functions? That sort of thing.

[Fznamznon]

clang/test/SemaSYCL/device-attributes.cpp

3 ↗

(On Diff #203785)

Actually there is no restrictions for adding this attribute to any function to outline device code so I just checked the simplest variant.

But I'm working on new patch which will put some requirements on function which is marked with sycl_kernel attribute.
This new patch will add generation of OpenCL kernel from function marked with sycl_kernel attribute. The main idea of this approach is described in this document (in this document generated kernel is called "kernel wrapper").
And to be able to generate OpenCL kernel using function marked with sycl_kernel attribute we put some requirements on this function, for example it must be a template function. You can find these requirements and example of proper function which can be marked with sycl_kernel in this comment .

[aaron.ballman]

clang/test/SemaSYCL/device-attributes.cpp

3 ↗

(On Diff #203785)

Actually there is no restrictions for adding this attribute to any function to outline device code so I just checked the simplest variant.

So there are no concerns about code like:

struct Base {
  __attribute__((sycl_kernel)) virtual void foo();
  virtual void bar();
};

struct Derived : Base {
  void foo() override;
  __attribute__((sycl_kernel)) void bar() override;
};

void f(Base *B, Derived *D) {
  // Will all of these "do the right thing"?
  B->foo();
  B->bar();

  D->foo();
  D->bar();
}

[Fznamznon]

Appled part of comments from @aaron.ballman:

• Fixed grammar and code style in all places except sycl_kernel docs
• Added a lit test which checks that sycl_device attribute implicitly added to proper declarations

[Fznamznon]

clang/include/clang/Basic/AttrDocs.td

263–264

I see. I will update documentation in the next version.

[bader]

clang/test/SemaSYCL/device-attributes.cpp

3 ↗

(On Diff #203785)

Actually there is no restrictions for adding this attribute to any function to outline device code so I just checked the simplest variant.
But I'm working on new patch which will put some requirements on function which is marked with sycl_kernel attribute.

@aaron.ballman, sorry for confusing. The usage scenarios should have been articulated more accurately.
We have only four uses of this attribute in our implementation:
https://github.com/intel/llvm/blob/sycl/sycl/include/CL/sycl/handler.hpp#L538 (lines 538-605).
All four uses are applied to member functions of cl::sycl::handler class and all of them have similar prototype (which is mentioned by Mariya in the previous comment:

namespace cl { namespace sycl {
class handler {
  template <typename KernelName, typename KernelType/*, ...*/>
  __attribute__((sycl_kernel)) void sycl_kernel_function(KernelType KernelFuncObj) {
    KernelFuncObj();
  }
};
}}

Here is the list of SYCL device compiler expectations with regard to the function marked with sycl_kernel attribute.

• Function template with at least one parameter is expected. The compiler generates OpenCL kernel and uses first template parameter as unique name to the generated OpenCL kernel. Host application uses this unique name to invoke the OpenCL kernel generated for the sycl_kernel_function specialized by this name and KernelType (which might be a lambda type).
• Function must have at least one parameter. First parameter expected to be a function object type (named or unnamed i.e. lambda). Compiler uses function object type field to generate OpenCL kernel parameters.

Aaron, I hope it makes more sense now.

We don't plan in any use cases other than in SYCL standard library implementation mentioned above.
If I understand you concerns correctly, you want to be sure that clang prohibits other uses of this attribute, which are not intended. Right?
What is the best way to do this? Add more negative tests cases and make sure that clang generate error diagnostic messages?

[keryell]

clang/test/SemaSYCL/device-attributes.cpp

3 ↗

(On Diff #203785)

If I understand you concerns correctly, you want to be sure that clang prohibits other uses of this attribute, which are not intended. Right?

But since it is an attribute to be used by SYCL run-time writers, I am not sure there is a lot of value in over-engineering the restrictions of its use. It diverts brain power from the real implementation & review and might even prevent innovation due to some creative use cases.

[Fznamznon]

Updated sycl_kernel attribute documentation.

[Fznamznon]

Fixed a couple coding style issues, renamed markDevice function with markSYCLDevice.

[aaron.ballman]

clang/test/SemaSYCL/device-attributes.cpp

3 ↗

(On Diff #203785)

If I understand you concerns correctly, you want to be sure that clang prohibits other uses of this attribute, which are not intended. Right?

Effectively, yes. I'd like to ensure that situations where the attribute does not do what the user expects are diagnosed. A good rule of thumb that I use is to diagnose (as a warning) situations where the attribute will be silently ignored, and diagnose (as an error) situations where applying the attribute would cause really bad results (like miscompiles, security concerns, etc).

What is the best way to do this? Add more negative tests cases and make sure that clang generate error diagnostic messages?

That's a good approach, yes. Though for the situations you describe, I'd probably just warn rather than err because it seems like it's harmless to ignore the attribute so long as the user knows it's being ignored.

3 ↗

(On Diff #203785)

But since it is an attribute to be used by SYCL run-time writers, I am not sure there is a lot of value in over-engineering the restrictions of its use. It diverts brain power from the real implementation & review and might even prevent innovation due to some creative use cases.

I disagree. Part of the real implementation is ensuring the attribute is not accidentally misused. It's frustrating for users to have an attribute silently ignored because it's easy to mistake that situation for the attribute behaving as expected.

[Fznamznon]

Added warning diagnostic for sycl_kernel attribute.

Now if the sycl_kernel attribute applied to a function which doesn't meet requirements for OpenCL kernel generation, attribute will be ignored and diagnostic will be emitted.

[Fznamznon]

Minor fix

[aaron.ballman]

clang/include/clang/Basic/AttrDocs.td

260

generate an OpenCL kernel

261

demonstrates the compiler's

278

defines the entry point

279

The compiler will

281

the compiler will add the "foo" function

282

More details about the compilation of functions for the device part can be found

284

of the code, the SYCL runtime

308

generate an OpenCL kernel

313

The function must be a template with at least two type template parameters.

314

generates an OpenCL kernel and uses the first template parameter as a unique name

315

The host application uses

318

The function must
The first parameter is required to be a

[aaron.ballman]

clang/include/clang/Basic/Attr.td

1074

Shouldn't this be FunctionTemplate instead?

clang/include/clang/Basic/AttrDocs.td

319

The compiler uses the function object type

321

The function must return void. The compiler reuses the body of marked functions to generate the OpenCL kernel body, and the OpenCL kernel must return void.

I'd move the "The sycl_kernel_function" sentence to its own paragraph rather than as part of the final bullet.

clang/include/clang/Basic/DiagnosticSemaKinds.td

10108–10109

I think this diagnostic should be split out into a few diagnostics that explicitly cover the requirements. Something like:
'sycl_kernel' attribute only applies to a %select{templated function|function returning 'void'|etc}0. It's best to avoid trying to send users to documentation if we can just tell them explicitly what they did wrong with their code.

clang/include/clang/Sema/Sema.h

11210 ↗

(On Diff #206873)

Can you add a const overload that returns a const container reference?

Also, why return the container rather than returning an iterator range from the container?

clang/lib/Sema/SemaDeclAttr.cpp

6417

Spurious newline above and missing a full stop at the end of the comment. Comments below are also missing full stops.

6418–6422

You can replace all this with a cast<FunctionDecl>(D) because the common attribute handler already verifies the subject is correct.

6424

I'd appreciate this being declared as a const pointer (same for the other nodes obtained through FT).

6427–6430

If you switch the subject to FunctionTemplate, then I believe this predicate can also go away.

clang/lib/Sema/SemaSYCL.cpp

65–66 ↗

(On Diff #206873)

I think there's some type confusion happening here. I would expect Elt to either be auto * or Func to be a const auto &. I suspect Elt should be declared as auto *.

[bader]

Hi @aaron.ballman,

Thanks a lot for the comments and sorry for the long delay. We've been working on complete implementation of the SYCL 1.2.1 specification.
Now I have more time to work on contributing the implementation to LLVM project.

I re-based the patch and started applying your suggestions.
In addition to that I'd like to investigate slightly different approach to outlining suggested by @ABataev at LLVM Dev. Meeting conference and utilize the infrastructure OpenMP compiler uses in CodeGen library to emit "device part" of the single source.

Thanks,
Alexey

[bader]

Applied comments from Aaron.

Two comments left unresolved:

• Split diagnostic message for sycl_kernel attribute into multiple messages. Will do tomorrow.
• Change attribute "subject" in TableGen file from "Function" to "FunctionTemplate". I need guidance here as I'm not sure how to make it work.

Refactored patch to re-use CodeGen infrastructure for emitting SYCL device code.
It turned out to be quite simple change - just two one-liner changes in ASTContext to say that only SYCL kernels must be emitted when we compile for SYCL device + similar change in the CodeGen to mark symbols which must be emitted.

Removed sycl_device attribute, which was required by previous implementation for device code outlining. I think we still might need this attribute to mark "non-kernel" symbols as "device code", so the compiler will emit even though they are not used by SYCL kernels. Anyway it's not required for device code outlining and shouldn't be part of this patch.

Enhanced CodeGen test to check that host part of the code is not emitted.

[bader]

clang/include/clang/Basic/Attr.td

1074

@aaron.ballman, I'm not sure.
I tried to use FunctionTemplate instead of Function, but I get following warning:

warning: 'sycl_kernel' attribute only applies to redeclarable templates

I investigated this a little and Sema passes Function declaration instead of FunctionTemplate to the function validating the attribute appertains to the right subject. I think it's because attributes are handled before FunctionTemplateDecl node is created.
Do we have an infrastructure to handle "FunctionTemplate" attributes?

I can't find any other attribute with FunctionTemplate subject to learn from...

[bader]

Applied two remaining comments from Aaron.

• Split diagnostics for sycl_kernel attribute to provide more informative message.
• Moved attribute target check to TableGen file. I stole a workaround for a function template subject emulation from @hfinkel C++ JIT compiler prototype (https://github.com/hfinkel/llvm-project-cxxjit/blob/cxxjit/clang/include/clang/Basic/Attr.td#L121).

[bader]

clang/test/SemaSYCL/device-attributes.cpp

35 ↗

(On Diff #228286)

Do we have to check each diagnostic message for both attribute spellings?

[bader]

@aaron.ballman, @Anastasia, could you take a look at new version of the patch, please?

[aaron.ballman]

clang/include/clang/Basic/DiagnosticSemaKinds.td

10118

Do you mean template function or function template? A function template is a template used to generate functions and a template function is a function produced by a template. I think you probably mean "function template" here.

10121–10122

This diagnostic reads a bit like you cannot do this: template <class N> when I think the actual restriction is that you cannot do this: template <int N>. Is that correct? If so, I think this could be worded as template parameter of a function template with the 'sycl_kernel' attribute must be a template type parameter.

Just double-checking, but you also intend to prohibit template template parameters? e.g., you can't do template <template <typename> typename C>

10124

Probably "function template" here as well.

10127

Same here.

[bader]

Applied code review comments.

[bader]

clang/include/clang/Basic/DiagnosticSemaKinds.td

10121–10122

This diagnostic reads a bit like you cannot do this: template <class N> when I think the actual restriction is that you cannot do this: template <int N>. Is that correct?

Yes. That is correct.

If so, I think this could be worded as template parameter of a function template with the 'sycl_kernel' attribute must be a template type parameter.

Thanks! Applied your wording.

Just double-checking, but you also intend to prohibit template template parameters? e.g., you can't do template <template <typename> typename C>

Currently we allow following use case: https://github.com/intel/llvm/blob/sycl/clang/test/SemaSYCL/mangle-kernel.cpp. I assume it qualifies as "template type" and not as "template template" parameter. Right?

Quoting SYCL specification $6.2 Naming of kernels (https://www.khronos.org/registry/SYCL/specs/sycl-1.2.1.pdf#page=250).

SYCL kernels are extracted from C++ source files and stored in an implementation- defined format. In the case of
the shared-source compilation model, the kernels have to be uniquely identified by both host and device compiler.
This is required in order for the host runtime to be able to load the kernel by using the OpenCL host runtime
interface.
From this requirement the following rules apply for naming the kernels:
• The kernel name is a C++ typename.
• The kernel needs to have a globally-visible name. In the case of a named function object type, the name can
be the typename of the function object, as long as it is globally-visible. In the case where it isn’t, a globally visible name has to be provided, as template parameter to the kernel invoking interface, as described in 4.8.5.
In C++11, lambdas do not have a globally-visible name, so a globally-visible typename has to be provided
in the kernel invoking interface, as described in 4.8.5.
• The kernel name has to be a unique identifier in the program.

We also have an extension, which lifts these restrictions/requirements when clang is used as host and device compiler. @erichkeane implemented built-in function (https://github.com/intel/llvm/pull/250) providing "unique identifier", which we use as a kernel name for lambda objects. But this is going to be a separate patch.

[bader]

clang/test/Misc/pragma-attribute-supported-attributes-list.test

134 ↗

(On Diff #230281)

It looks like this change is not needed anymore. This check fails on my machine with the latest version of the patch.

@aaron.ballman, I'm not sure if this is a problem of the implementation or test issue.
Do I understand correctly that this test validates the list of the attributes which can be applied using #pragma clang?
If so, removing this check seems to be okay. We need only [[clang::sycl_kernel]] or __attribute__((sycl_kernel)) to work.

[aaron.ballman]

clang/include/clang/Basic/DiagnosticSemaKinds.td

10121–10122

Currently we allow following use case: https://github.com/intel/llvm/blob/sycl/clang/test/SemaSYCL/mangle-kernel.cpp. I assume it qualifies as "template type" and not as "template template" parameter. Right?

Yeah, those are template types. A template template parameter would be: https://godbolt.org/z/9kwbW9
In that example, C is a template template parameter and Ty is a template type parameter. The part I'm a bit unclear on is why a template template parameter should be disallowed (I believe it names a type, as opposed to a non-type template parameter which names a value)?

clang/test/Misc/pragma-attribute-supported-attributes-list.test

134 ↗

(On Diff #230281)

Your understanding is correct, and I think it's a bug if you don't need to add an entry here for SYCLKernel. @arphaman, WDYT?

[bader]

Applied code review comments from Aaron.

Allow template template parameters for function templates marked with sycl_kernel attribute.

[bader]

clang/include/clang/Basic/DiagnosticSemaKinds.td

10121–10122

I think Mariya implemented this restriction to avoid usages not required for SYCL kernel support implementation in run-time library. As it was mentioned before, this attribute is intended to be used by SYCL run-time library only and current implantation do not require template template parameter support.
I think that this might be useful for alternative implementations, so I updated the patch to restrict non-type template parameters only.

clang/test/Misc/pragma-attribute-supported-attributes-list.test

134 ↗

(On Diff #230281)

I turned out that the workaround I added to allow only function templates affected this test (described in this comment https://reviews.llvm.org/D60455#1742083).

I.e.
def FunctionTmpl

: SubsetSubject<Function, [{S->getTemplatedKind() ==
                             FunctionDecl::TK_FunctionTemplate}],
                "function templates">;

I also noted that there is no check for artificial attribute which uses the same approach to limit the subject to "inline functions".

https://github.com/llvm/llvm-project/blob/master/clang/include/clang/Basic/Attr.td#L652
https://github.com/llvm/llvm-project/blob/master/clang/include/clang/Basic/Attr.td#L122

[bader]

Ping.

[ABataev]

clang/lib/CodeGen/CodeGenModule.cpp

2477 ↗

(On Diff #230310)

Need to check if the decl must be emitted at all.

[bader]

clang/lib/CodeGen/CodeGenModule.cpp

2477 ↗

(On Diff #230310)

Let me check that I get it right. You suggest adding if (MustBeEmitted(Global)), right?

if (LangOpts.SYCLIsDevice && Global->hasAttr<SYCLKernelAttr>() && MustBeEmitted(Global)) {
  ...
  addDeferredDeclToEmit(GD);
  return;
}

[ABataev]

clang/lib/CodeGen/CodeGenModule.cpp

2477 ↗

(On Diff #230310)

Yes

[bader]

clang/lib/CodeGen/CodeGenModule.cpp

2477 ↗

(On Diff #230310)

Okay. Making this change requires additional adjustments in the patch and I have a few options.
In this patch we do not add any logic forcing compiler to emit SYCL kernel. This logic is supposed to be added by follow-up patch (currently under SYCL working group review here https://github.com/intel/llvm/pull/249), which add code emitting "externally visible" OpenCL kernel calling function object passed to SYCL kernel function.

I can:

1. Temporally remove CodeGen test and add updated version back with the follow-up patch
2. Do change making SYCL kernels "externally visible" and revert this change with the follow-up patch (this is kind of current logic which emits SYCL kernels unconditionally)
3. Merge two patches and submit them together, but I assume it will significantly increase the size of the patch.

[ABataev]

clang/lib/CodeGen/CodeGenModule.cpp

2477 ↗

(On Diff #230310)

Probably, better would be to split the patch

[aaron.ballman]

The attribute bits LGTM aside from a wording nit with the diagnostic; I have no opinion on the CodeGen question.

clang/include/clang/Basic/DiagnosticSemaKinds.td

10122

can't -> cannot

[bader]

Applied code review suggestions.

• Split the patch into two parts. This patch contains only Sema part and LLVM IR generation part will be added separately. Updated commit message.
• Replace can't with cannot.

[bader]

Fixed typo in the commit message: complier -> compiler.

[Anastasia]

Sorry, I don't have capacity currently to review this and I don't want to be blocking it either.

[bader]

In D60455#1762804, @Anastasia wrote:

Sorry, I don't have capacity currently to review this and I don't want to be blocking it either.

@Anastasia, thanks for finding time for reviewing previous revisions of the patch. I really appreciate your comments.

[bader]

Minor update adjusting to the recent changes.

Updated comment "The 'sycl_kernel' attribute applies only to functions" -> "The 'sycl_kernel' attribute applies only to function templates".
Renamed tests from "device-attributes*" to "kernel-attribute*".

[Fznamznon]

A couple of minor comments.

clang/include/clang/Basic/AttrDocs.td

313

@bader , could you please apply this too?

314

I'm not an expert in English, so you can ignore it if I'm wrong, but a phrase like "uses parameter as a name to the kernel" seems strange. Maybe "for kernel"?

317

(which might be a lambda or a function object type).

[bader]

Applied comments from @Fznamznon.

[aaron.ballman]

LGTM with some testing requests.

clang/test/SemaSYCL/kernel-attribute.cpp

5–6

Still missing a test that the attribute is ignored when SYCL is not enabled.

7

This test should be on a templated function (we already demonstrated it only applies to templated functions, so the check for the argument is not what is failing).

8

Same here.

[bader]

clang/test/SemaSYCL/kernel-attribute.cpp

5–6

Still missing a test that the attribute is ignored when SYCL is not enabled.

I think clang/test/SemaSYCL/kernel-attribute-on-non-sycl.cpp should check that. Please, let me know if you mean something else.

This test should be on a templated function (we already demonstrated it only applies to templated functions, so the check for the argument is not what is failing).

Nice catch. Thanks!

[bader]

Applied @aaron.ballman suggestions to kernel-attribute.cpp test

[Fznamznon]

LGTM with a couple of minor comments.

clang/include/clang/Basic/AttrDocs.td

273

Sorry for late catch, but there is a little bug in this SYCL code: index is one-dimensional id, so calling subscript operator with any value other than 0 is a bug.

319

There are two spaces between "." and "The" at the end of line 319.

[aaron.ballman]

clang/test/SemaSYCL/kernel-attribute.cpp

5–6

I think clang/test/SemaSYCL/kernel-attribute-on-non-sycl.cpp should check that. Please, let me know if you mean something else.

Oh, you're correct, that was the test I was hoping for!

[bader]

Fixed SYCL code example for sycl_kernel attribute documentation and commit message.

[bader]

I hope all comments from are @Fznamznon and @aaron.ballman are applied.
@ABataev, do you have any other comments?

clang/include/clang/Basic/Attr.td

def SharedVar : SubsetSubject<Var,
                               "global variables">;
 
 def GlobalVar : SubsetSubject<Var,
                              [{S->hasGlobalStorage()}], "global variables">;
 
 def InlineFunction : SubsetSubject<Function,
                              [{S->isInlineSpecified()}], "inline functions">;
 
+def FunctionTmpl
+    : SubsetSubject<Function, [{S->getTemplatedKind() ==
+                                 FunctionDecl::TK_FunctionTemplate}],
+                    "function templates">;
+
 // FIXME: this hack is needed because DeclNodes.td defines the base Decl node
 // type to be a class, not a definition. This makes it impossible to create an
 // attribute subject which accepts a Decl. Normally, this is not a problem,
 // because the attribute can have no Subjects clause to accomplish this. But in
 // the case of a SubsetSubject, there's no way to express it without this hack.
 def DeclBase : AttrSubject;
 def FunctionLike : SubsetSubject<DeclBase,
                                  [{S->getFunctionType(false) != nullptr}],
class LangOpt<string name, code customCode = [{}]> {
   // A custom predicate, written as an expression evaluated in a context with
   // "LangOpts" bound.
   code CustomCode = customCode;
 }
 def MicrosoftExt : LangOpt<"MicrosoftExt">;
 def Borland : LangOpt<"Borland">;
 def CUDA : LangOpt<"CUDA">;
 def HIP : LangOpt<"HIP">;
+def SYCL : LangOpt<"SYCLIsDevice">;
 def COnly : LangOpt<"COnly", "!LangOpts.CPlusPlus">;
 def CPlusPlus : LangOpt<"CPlusPlus">;
 def OpenCL : LangOpt<"OpenCL">;
 def RenderScript : LangOpt<"RenderScript">;
 def ObjC : LangOpt<"ObjC">;
 def BlocksSupported : LangOpt<"Blocks">;
 def ObjCAutoRefCount : LangOpt<"ObjCAutoRefCount">;
 def ObjCNonFragileRuntime : LangOpt<"ObjCNonFragileRuntime",

 
 def CUDAShared : InheritableAttr {
   let Spellings = [GNU<"shared">, Declspec<"__shared__">];
   let Subjects = SubjectList<[Var]>;
   let LangOpts = [CUDA];
   let Documentation = [Undocumented];
 }
 
+def SYCLKernel : InheritableAttr {
+  let Spellings = [Clang<"sycl_kernel">];
+  let Subjects = SubjectList<[FunctionTmpl]>;
+  let LangOpts = [SYCL];
+  let Documentation = [SYCLKernelDocs];
+}
+
 def C11NoReturn : InheritableAttr {
   let Spellings = [Keyword<"_Noreturn">];
   let Subjects = SubjectList<[Function], ErrorDiag>;
   let SemaHandler = 0;
   let Documentation = [C11NoReturnDocs];
 }
 
 def CXX11NoReturn : InheritableAttr {
   let Spellings = [CXX11<"", "noreturn", 200809>];
   let Subjects = SubjectList<[Function], ErrorDiag>;
   let Documentation = [CXX11NoReturnDocs];
 }
 

clang/include/clang/Basic/AttrDocs.td

 
 It is also possible to specify a CPU name of ``generic`` which will be resolved
 if the executing processor doesn't satisfy the features required in the CPU
 name. The behavior of a program executing on a processor that doesn't satisfy
 any option of a multiversioned function is undefined.
   }];
 }
 
+def SYCLKernelDocs : Documentation {
+  let Category = DocCatFunction;
+  let Content = [{
+The ``sycl_kernel`` attribute specifies that a function template will be used
+to outline device code and to generate an OpenCL kernel.
+Here is a code example of the SYCL program, which demonstrates the compiler's
+outlining job:
+.. code-block:: c++
+
+  int foo(int x) { return ++x; }
+
+  using namespace cl::sycl;
+  queue Q;
+  buffer<int, 1> a(range<1>{1024});
+  Q.submit([&](handler& cgh) {
+    auto A = a.get_access<access::mode::write>(cgh);
+    cgh.parallel_for<init_a>(range<1>{1024}, [=](id<1> index) {
+      A[index] = index[0] + foo(42);
+    });
+  }
+
+A C++ function object passed to the ``parallel_for`` is called a "SYCL kernel".
+A SYCL kernel defines the entry point to the "device part" of the code. The
+compiler will emit all symbols accessible from a "kernel". In this code
+example, the compiler will emit "foo" function.  More details about the
+compilation of functions for the device part can be found in the SYCL 1.2.1
+specification Section 6.4.
+To show to the compiler entry point to the "device part" of the code, the SYCL
+runtime can use the ``sycl_kernel`` attribute in the following way:
+.. code-block:: c++
+namespace cl {
+namespace sycl {
+class handler {
+  template <typename KernelName, typename KernelType/*, ...*/>
+  __attribute__((sycl_kernel)) void sycl_kernel_function(KernelType KernelFuncObj) {
+    // ...
+    KernelFuncObj();
+  }
+
+  template <typename KernelName, typename KernelType, int Dims>
+  void parallel_for(range<Dims> NumWorkItems, KernelType KernelFunc) {
+#ifdef __SYCL_DEVICE_ONLY__
+    sycl_kernel_function<KernelName, KernelType, Dims>(KernelFunc);
+#else
+    // Host implementation
+#endif
+  }
+};
+} // namespace sycl
+} // namespace cl
+
+The compiler will also generate an OpenCL kernel using the function marked with
+the ``sycl_kernel`` attribute.
+Here is the list of SYCL device compiler expectations with regard to the
+function marked with the ``sycl_kernel`` attribute:
+
+- The function must be a template with at least two type template parameters.
+  The compiler generates an OpenCL kernel and uses the first template parameter
+  as a unique name for the generated OpenCL kernel. The host application uses
+  this unique name to invoke the OpenCL kernel generated for the SYCL kernel
+  specialized by this name and second template parameter ``KernelType`` (which
+  might be an unnamed function object type).
+- The function must have at least one parameter. The first parameter is
+  required to be a function object type (named or unnamed i.e. lambda). The
+  compiler uses function object type fields to generate OpenCL kernel
+  parameters.
+- The function must return void. The compiler reuses the body of marked functions to
+  generate the OpenCL kernel body, and the OpenCL kernel must return `void`.
+
+The SYCL kernel in the previous code sample meets these expectations.
+  }];
+}
+
 def C11NoReturnDocs : Documentation {
   let Category = DocCatFunction;
   let Content = [{
 A function declared as ``_Noreturn`` shall not return to its caller. The
 compiler will generate a diagnostic for a function declared as ``_Noreturn``
 that appears to be capable of returning to its caller. Despite being a type
 specifier, the ``_Noreturn`` attribute cannot be specified on a function
 pointer type.

clang/include/clang/Basic/DiagnosticSemaKinds.td

def warn_noderef_to_dereferenceable_pointer : Warning<
   "casting to dereferenceable pointer removes 'noderef' attribute">, InGroup<NoDeref>;
 
 def err_builtin_launder_invalid_arg : Error<
   "%select{non-pointer|function pointer|void pointer}0 argument to "
   "'__builtin_launder' is not allowed">;
 
 def err_preserve_field_info_not_field : Error<
   "__builtin_preserve_field_info argument %0 not a field access">;
 def err_preserve_field_info_not_const: Error<
   "__builtin_preserve_field_info argument %0 not a constant">;
 
 def err_bit_cast_non_trivially_copyable : Error<
   "__builtin_bit_cast %select{source|destination}0 type must be trivially copyable">;
 def err_bit_cast_type_size_mismatch : Error<
   "__builtin_bit_cast source size does not equal destination size (%0 vs %1)">;
+
+// SYCL-specific diagnostics
+def warn_sycl_kernel_num_of_template_params : Warning<
+  "'sycl_kernel' attribute only applies to a function template with at least"
+  " two template parameters">, InGroup<IgnoredAttributes>;
+def warn_sycl_kernel_invalid_template_param_type : Warning<
+  "template parameter of a function template with the 'sycl_kernel' attribute"
+  " cannot be a non-type template parameter">, InGroup<IgnoredAttributes>;
+def warn_sycl_kernel_num_of_function_params : Warning<
+  "function template with 'sycl_kernel' attribute must have a single parameter">,
+  InGroup<IgnoredAttributes>;
+def warn_sycl_kernel_return_type : Warning<
+  "function template with 'sycl_kernel' attribute must have a 'void' return type">,
+  InGroup<IgnoredAttributes>;
+
 } // end of sema component.

clang/lib/Sema/SemaDeclAttr.cpp

if (AL.getAttrName()->getName().find("read_write") != StringRef::npos) {
         return;
       }
     }
   }
 
   D->addAttr(::new (S.Context) OpenCLAccessAttr(S.Context, AL));
 }
 
+static void handleSYCLKernelAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
+  // The 'sycl_kernel' attribute applies only to function templates.
+  const auto *FD = cast<FunctionDecl>(D);
+  const FunctionTemplateDecl *FT = FD->getDescribedFunctionTemplate();
+  assert(FT && "Function template is expected");
+
+  // Function template must have at least two template parameters.
+  const TemplateParameterList *TL = FT->getTemplateParameters();
+  if (TL->size() < 2) {
+    S.Diag(FT->getLocation(), diag::warn_sycl_kernel_num_of_template_params);
+    return;
+  }
+
+  // Template parameters must be typenames.
+  for (unsigned I = 0; I < 2; ++I) {
+    const NamedDecl *TParam = TL->getParam(I);
+    if (isa<NonTypeTemplateParmDecl>(TParam)) {
+      S.Diag(FT->getLocation(),
+             diag::warn_sycl_kernel_invalid_template_param_type);
+      return;
+    }
+  }
+
+  // Function must have at least one argument.
+  if (getFunctionOrMethodNumParams(D) != 1) {
+    S.Diag(FT->getLocation(), diag::warn_sycl_kernel_num_of_function_params);
+    return;
+  }
+
+  // Function must return void.
+  QualType RetTy = getFunctionOrMethodResultType(D);
+  if (!RetTy->isVoidType()) {
+    S.Diag(FT->getLocation(), diag::warn_sycl_kernel_return_type);
+    return;
+  }
+
+  handleSimpleAttribute<SYCLKernelAttr>(S, D, AL);
+}
+
 static void handleDestroyAttr(Sema &S, Decl *D, const ParsedAttr &A) {
   if (!cast<VarDecl>(D)->hasGlobalStorage()) {
     S.Diag(D->getLocation(), diag::err_destroy_attr_on_non_static_var)
         << (A.getKind() == ParsedAttr::AT_AlwaysDestroy);
     return;
   }
 
   if (A.getKind() == ParsedAttr::AT_AlwaysDestroy)
case ParsedAttr::AT_FlagEnum:
     handleSimpleAttribute<FlagEnumAttr>(S, D, AL);
     break;
   case ParsedAttr::AT_EnumExtensibility:
     handleEnumExtensibilityAttr(S, D, AL);
     break;
   case ParsedAttr::AT_Flatten:
     handleSimpleAttribute<FlattenAttr>(S, D, AL);
     break;
+  case ParsedAttr::AT_SYCLKernel:
+    handleSYCLKernelAttr(S, D, AL);
+    break;
   case ParsedAttr::AT_Format:
     handleFormatAttr(S, D, AL);
     break;
   case ParsedAttr::AT_FormatArg:
     handleFormatArgAttr(S, D, AL);
     break;
   case ParsedAttr::AT_Callback:
     handleCallbackAttr(S, D, AL);

clang/test/SemaSYCL/kernel-attribute-on-non-sycl.cpp

+// RUN: %clang_cc1 -std=c++11 -fsyntax-only -fsycl-is-device -verify %s
+// RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify -x c++ %s
+
+#ifndef __SYCL_DEVICE_ONLY__
+// expected-warning@+7 {{'sycl_kernel' attribute ignored}}
+// expected-warning@+8 {{'sycl_kernel' attribute ignored}}
+#else
+// expected-no-diagnostics
+#endif
+
+template <typename T, typename A, int B>
+__attribute__((sycl_kernel)) void foo(T P);
+template <typename T, typename A, int B>
+[[clang::sycl_kernel]] void foo1(T P);

clang/test/SemaSYCL/kernel-attribute.cpp

+// RUN: %clang_cc1 -std=c++11 -fsyntax-only -fsycl-is-device -verify %s
+
+// Only function templates
+[[clang::sycl_kernel]] int gv2 = 0; // expected-warning {{'sycl_kernel' attribute only applies to function templates}}
+__attribute__((sycl_kernel)) int gv3 = 0; // expected-warning {{'sycl_kernel' attribute only applies to function templates}}
+
+__attribute__((sycl_kernel)) void foo(); // expected-warning {{'sycl_kernel' attribute only applies to function templates}}
+[[clang::sycl_kernel]] void foo1(); // expected-warning {{'sycl_kernel' attribute only applies to function templates}}
+
+// Attribute takes no arguments
+template <typename T, typename A>
+__attribute__((sycl_kernel(1))) void foo(T P); // expected-error {{'sycl_kernel' attribute takes no arguments}}
+template <typename T, typename A, int I>
+[[clang::sycl_kernel(1)]] void foo1(T P);// expected-error {{'sycl_kernel' attribute takes no arguments}}
+
+// At least two template parameters
+template <typename T>
+__attribute__((sycl_kernel)) void foo(T P); // expected-warning {{'sycl_kernel' attribute only applies to a function template with at least two template parameters}}
+template <typename T>
+[[clang::sycl_kernel]] void foo1(T P); // expected-warning {{'sycl_kernel' attribute only applies to a function template with at least two template parameters}}
+
+// First two template parameters cannot be non-type template parameters
+template <typename T, int A>
+__attribute__((sycl_kernel)) void foo(T P); // expected-warning {{template parameter of a function template with the 'sycl_kernel' attribute cannot be a non-type template parameter}}
+template <int A, typename T>
+[[clang::sycl_kernel]] void foo1(T P); // expected-warning {{template parameter of a function template with the 'sycl_kernel' attribute cannot be a non-type template parameter}}
+
+// Must return void
+template <typename T, typename A>
+__attribute__((sycl_kernel)) int foo(T P); // expected-warning {{function template with 'sycl_kernel' attribute must have a 'void' return type}}
+template <typename T, typename A>
+[[clang::sycl_kernel]] int foo1(T P); // expected-warning {{function template with 'sycl_kernel' attribute must have a 'void' return type}}
+
+// Must take at least one argument
+template <typename T, typename A>
+__attribute__((sycl_kernel)) void foo(); // expected-warning {{function template with 'sycl_kernel' attribute must have a single parameter}}
+template <typename T, typename A>
+[[clang::sycl_kernel]] void foo1(T t, A a); // expected-warning {{function template with 'sycl_kernel' attribute must have a single parameter}}
+
+// No diagnostics
+template <typename T, typename A>
+__attribute__((sycl_kernel)) void foo(T P);
+template <typename T, typename A, int I>
+[[clang::sycl_kernel]] void foo1(T P);