keryell added a comment.

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

  rG LLVM Github Monorepo


cfe-commits mailing list

Reply via email to