bader added a comment.
To give more context to the question, I'd like to clarify the use case of new
attributes.
As @Fznamznon already mentioned in previous comments SYCL is not supposed to
expose any non-standard extensions to a user.
Here is code example of the SYCL program, which demonstrate the need for these
attributes:
C++
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 compiler need to compile lambda function passed to
`cl::sycl::handler::parallel_for` method and function `foo` called from this
lambda function.
NOTE: compiler must ignore `bar` function when we "device" part of the single
source code.
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:
https://github.com/intel/llvm/blob/sycl/sycl/include/CL/sycl/handler.hpp#L267
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.
What we are looking for is whether we should add SYCL specific code or we can
come up with something more generic to avoid logic/code duplication with
already existing functionality.
BTW: Mariya, I think we might need to use `sycl_device` attribute to mark
functions, which are called from the different translation units, i.e. compiler
can't identify it w/o user's help.
SYCL specification proposes to use special macro as "device function marker",
but I guess we can have additional "spellings" in the clang.
NOTE2: @Anastasia, https://reviews.llvm.org/D60454 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.
Repository:
rG LLVM Github Monorepo
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D60455/new/
https://reviews.llvm.org/D60455
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
