[all-commits] [llvm/llvm-project] c094e7: [SYCL] Add sycl_kernel attribute for accelerated c...

[Mariya Podchishchaeva via All-commits]

  Branch: refs/heads/master
  Home:   https://github.com/llvm/llvm-project
  Commit: c094e7dc4b3f9d1c1e590b008bb1cc46e3496abd
      https://github.com/llvm/llvm-project/commit/c094e7dc4b3f9d1c1e590b008bb1cc46e3496abd
  Author: Mariya Podchishchaeva <mariya.podchishchaeva at intel.com>
  Date:   2019-12-03 (Tue, 03 Dec 2019)

  Changed paths:
    M clang/include/clang/Basic/Attr.td
    M clang/include/clang/Basic/AttrDocs.td
    M clang/include/clang/Basic/DiagnosticSemaKinds.td
    M clang/lib/Sema/SemaDeclAttr.cpp
    A clang/test/SemaSYCL/kernel-attribute-on-non-sycl.cpp
    A clang/test/SemaSYCL/kernel-attribute.cpp

  Log Message:
  -----------
  [SYCL] Add sycl_kernel attribute for accelerated code outlining

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] + 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.

Reviewers: jlebar, keryell, Naghasan, ABataev, Anastasia, bader, aaron.ballman, rjmccall, rsmith

Reviewed By: keryell, bader

Subscribers: mgorny, OlegM, ArturGainullin, agozillon, aaron.ballman, ebevhan, Anastasia, cfe-commits

Tags: #clang

Differential Revision: https://reviews.llvm.org/D60455

Signed-off-by: Alexey Bader <alexey.bader at intel.com>