This is an archive of the discontinued LLVM Phabricator instance.

[SYCL] Add sycl_kernel attribute for accelerated code outlining
ClosedPublic

Authored by bader on Apr 9 2019, 5:36 AM.

Details

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.

Diff Detail

Event Timeline

There are a very large number of changes, so older changes are hidden. Show Older Changes
bader added inline comments.May 27 2019, 7:16 AM
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.

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 added inline comments.May 28 2019, 4:35 AM
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 updated this revision to Diff 201641.May 28 2019, 5:07 AM

Applied comments from @Anastasia

  • Added documentation for sycl_kernel function
  • Added comments to Sema.h
  • Added -std=c++11 to test run lines
Anastasia added inline comments.May 30 2019, 10:53 AM
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 added inline comments.May 31 2019, 5:52 AM
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 added inline comments.Jun 3 2019, 3:00 AM
clang/test/CodeGenSYCL/device-functions.cpp
24 ↗(On Diff #201641)

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

Fznamznon added inline comments.Jun 3 2019, 3:28 AM
clang/test/CodeGenSYCL/device-functions.cpp
24 ↗(On Diff #201641)

There will be no SPIR calling convention for device functions.

Anastasia added inline comments.Jun 3 2019, 7:54 AM
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 updated this revision to Diff 203785.Jun 10 2019, 2:24 AM

Applied comments from @Anastasia

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

@Anastasia, do you have additional comments?

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

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 added inline comments.Jun 18 2019, 8:01 AM
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 added inline comments.Jun 18 2019, 3:15 PM
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 updated this revision to Diff 205663.Jun 19 2019, 1:42 PM

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 added inline comments.Jun 19 2019, 1:45 PM
clang/include/clang/Basic/AttrDocs.td
263–264

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

bader added inline comments.Jun 19 2019, 2:00 PM
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 added inline comments.Jun 19 2019, 6:55 PM
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 this revision to Diff 205813.Jun 20 2019, 7:54 AM

Updated sycl_kernel attribute documentation.

Fznamznon updated this revision to Diff 205831.Jun 20 2019, 8:58 AM

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

aaron.ballman added inline comments.Jun 24 2019, 1:28 PM
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 updated this revision to Diff 206861.Jun 27 2019, 7:17 AM

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.

aaron.ballman added inline comments.Jul 1 2019, 8:13 AM
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 added inline comments.Jul 1 2019, 8:13 AM
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
10104–10105

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 added a comment.Nov 7 2019, 3:34 AM

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 updated this revision to Diff 228286.Nov 7 2019, 11:53 AM

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 added inline comments.Nov 7 2019, 11:55 AM
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 commandeered this revision.Nov 7 2019, 12:01 PM
bader edited reviewers, added: Fznamznon; removed: bader.
bader updated this revision to Diff 228868.Nov 12 2019, 4:37 AM
bader added a subscriber: hfinkel.

Applied two remaining comments from Aaron.

bader marked an inline comment as done.Nov 12 2019, 4:38 AM
bader added inline comments.
clang/test/SemaSYCL/device-attributes.cpp
35 ↗(On Diff #228286)

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

bader added a comment.Nov 19 2019, 6:00 AM

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

aaron.ballman added inline comments.Nov 20 2019, 6:37 AM
clang/include/clang/Basic/DiagnosticSemaKinds.td
10114

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.

10117–10118

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>

10120

Probably "function template" here as well.

10123

Same here.

bader updated this revision to Diff 230281.Nov 20 2019, 9:23 AM
bader marked 5 inline comments as done.

Applied code review comments.

bader added inline comments.
clang/include/clang/Basic/DiagnosticSemaKinds.td
10117–10118

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 marked an inline comment as done.Nov 20 2019, 9:54 AM
bader added inline comments.
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 added a subscriber: arphaman.
aaron.ballman added inline comments.
clang/include/clang/Basic/DiagnosticSemaKinds.td
10117–10118

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 updated this revision to Diff 230310.Nov 20 2019, 12:18 PM

Applied code review comments from Aaron.

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

bader marked 3 inline comments as done.Nov 20 2019, 12:33 PM
bader added inline comments.
clang/include/clang/Basic/DiagnosticSemaKinds.td
10117–10118

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

ABataev added inline comments.Nov 27 2019, 6:23 AM
clang/lib/CodeGen/CodeGenModule.cpp
2477 ↗(On Diff #230310)

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

bader marked an inline comment as done.Nov 27 2019, 7:25 AM
bader added inline comments.
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 added inline comments.Nov 27 2019, 7:32 AM
clang/lib/CodeGen/CodeGenModule.cpp
2477 ↗(On Diff #230310)

Yes

bader marked an inline comment as done.Nov 27 2019, 9:21 AM
bader added inline comments.
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 added inline comments.Nov 27 2019, 11:00 AM
clang/lib/CodeGen/CodeGenModule.cpp
2477 ↗(On Diff #230310)

Probably, better would be to split the patch

aaron.ballman accepted this revision.Nov 27 2019, 1:40 PM

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
10118

can't -> cannot

bader updated this revision to Diff 231386.Nov 28 2019, 2:33 AM

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 retitled this revision from [SYCL] Implement SYCL device code outlining to [SYCL] Add sycl_kernel attribute for accelerated code outlining.Nov 28 2019, 3:04 AM
bader edited the summary of this revision. (Show Details)
bader updated this revision to Diff 231396.Nov 28 2019, 3:05 AM

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

Anastasia resigned from this revision.Nov 28 2019, 3:23 AM

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

This revision is now accepted and ready to land.Nov 28 2019, 3:23 AM
bader added a comment.Nov 28 2019, 3:36 AM

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 updated this revision to Diff 231506.Nov 29 2019, 2:43 AM

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*".

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 updated this revision to Diff 231595.Nov 30 2019, 12:37 PM
bader marked 4 inline comments as done.

Applied comments from @Fznamznon.

aaron.ballman accepted this revision.Dec 1 2019, 8:40 AM

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 marked 5 inline comments as done.Dec 2 2019, 12:21 AM
bader added inline comments.
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 updated this revision to Diff 231644.Dec 2 2019, 12:22 AM
bader marked an inline comment as done.

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

Fznamznon accepted this revision.Dec 2 2019, 4:20 AM

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 added inline comments.Dec 2 2019, 4:44 AM
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 updated this revision to Diff 231679.Dec 2 2019, 5:01 AM
bader marked 2 inline comments as done.

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

bader added a comment.Dec 2 2019, 5:03 AM

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

This revision was automatically updated to reflect the committed changes.