[PATCH] D60455: [SYCL] Add support for SYCL device attributes

[Ronan Keryell via Phabricator via cfe-commits]

keryell added a comment.

In D60455#1470402 <https://reviews.llvm.org/D60455#1470402>, @Anastasia wrote:

>




> In the interest of speeding up the upstreaming work, would you be able to highlight the differences and similarity at least for SYCL and OpenCL kernel modes? Not sure if you are familiar enough with both. Because apart from the public announcements I can't see what has been changed in SYCL that would disallow to use OpenCL mode. It would be a valuable input to determine the implementation choices.

SYCL is similar to OpenMP 5 for C++, where you use only C++ classes instead of `#pragma`. So it is quite C++-friendlier than OpenMP.
But that means also there is not the same concept of explicit kernel like in OpenCL or CUDA. In OpenCL or CUDA, when there is a function with a specific attribute, you know it is a kernel and you can compile as such.

In SYCL or OpenMP you need an outliner that will estimate what should be executed as an heterogeneous kernel, split the code between the host side and the device side, add some glue/stub to implement an RPC between the host and the device, manage potentially some allocation/memory transfers, etc. This is quite more complex than compiling OpenCL, CUDA or other graphics shader languages. This is also why, while SYCL is technically pure standard C++, you need some specific compiler magic to do the code massaging to have everything working well between a host and some devices.

The attribute we discuss here is just an implementation detail to help the coordination between the compiler and the SYCL frontend classes to mark some area to outline, without relying to do some precise pattern matching, allowing more flexibility in the runtime without changing the compiler every time. So while it defines a zone to be outlined as a kernel, it is not really a kernel in the sense of OpenCL.

In triSYCL I made some completely different choices, using late outlining in LLVM and detecting some specific functions such as `cl::sycl::detail::instantiate_kernel<KernelName>()` that defines some stuff I want to outline https://github.com/triSYCL/triSYCL/blob/master/doc/architecture.rst#low-level-view-of-the-device-compiler-workflow
For me an attribute was not an option because I wanted to change Clang as little as possible. But at the end, I think it is quite more brittle than doing early outlining in Clang as discussed here, which also requires quite more knowledge of Clang than I have. :-)

So at the end, I think we should use a different keyword from OpenCL or CUDA because the semantics is different.


Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D60455/new/

https://reviews.llvm.org/D60455