Anastasia added a comment.

> Our current approach is to add an attribute, which SYCL runtime will use to 
> mark code passed to `cl::sycl::handler::parallel_for` as "kernel functions". 
> Obviously runtime library can't mark `foo` as "device" code - this is a 
> compiler job: to traverse all symbols accessible from kernel functions and 
> add them to the "device part" of the code.
> Here is a link to the code in the SYCL runtime using `sycl_kernel` attribute: 
> I'm quite sure something similar should happen for other "single source" 
> programming models like OpenMP/CUDA, except these attributes are exposed to 
> the user and there is a specific requirement on attributes/pragma/keyword 
> names.

Just to understand how this will work. I would imagine you can have a device 
function definition preceding its use. When the function is being parsed it's 
not known yet whether it will be called from the device or not, so it won't be 
possible to set the language mode correctly and hence provide the right 
diagnostics. So is the plan to launch a separate parsing phase then just to 
extract the call graph and annotate the device functions?

> NOTE2: @Anastasia, makes impossible to 
> replace `sycl_kernel` attribute with `__kernel` attribute. I mean we still 
> can enable it for SYCL extension, but we will need SYCL specific 
> customization in this case as we apply `kernel` attribute to a template 
> function.

Ok, my question is whether you are planning to duplicate the same logic as for 
OpenCL kernel which doesn't really seem like an ideal design choice. Is this 
the only difference then we can simply add an extra check for SYCL compilation 
mode in this template handling case. The overall interaction between OpenCL and 
SYCL implementation is still a very big unknown to me so it's not very easy to 
judge about the implementations details...

  rG LLVM Github Monorepo


cfe-commits mailing list

Reply via email to