Page MenuHomePhabricator

[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
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
318

The function must
The first parameter is required to be a

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

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
10122

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.