Fznamznon added a comment.
I tried to reuse OpenCL kernel attribute with "__kernel" keyword in our current
SYCL implementation. PR with this try is here -
https://github.com/intel/llvm/pull/97
Now It looks feasible but with a couple notes:
From SYCL specification "SYCL is designed to be as close to standard C++ as
possible. Standard C++ compiler can compile the SYCL programs and they will run
correctly on host CPU." So SYCL doesn't provide non-standard `kernel` keyword
which is provided by OpenCL. Due this fact it's not possible to add `kernel`
keyword as in OpenCL, it will prevent compilation of following valid SYCL code:
int foo(int kernel) { return ++kernel; } // If "kernel" will be a keyword
like in OpenCL, here will be a error
…
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);
});
}
...
So I added only `__kernel` keyword for SYCL because in C++ identifiers which
start with `__` are reserved for compiler internals.
Next note:
In our current implementation actually not quite that function which is marked
with `sycl_kernel` (or `__kernel`, whatever) will be real OpenCL kernel in
produced module. In SYCL all shared between host and device memory objects
(buffers/images, these objects map to OpenCL buffers and images) can be
accessed through special `accessor` classes. SYCL also has special mechanism
for passing kernel arguments from host to device, if in OpenCL you need to do
`clSetKernelArg`, in SYCL all kernel arguments are captures/fields of
lambda/functor which is passed to `parallel_for` (See code snippet above, here
one kernel argument - accessor `A` ). To map to OpenCL setting kernel arguments
mechanism we added generation of some "kernel wrapper" function inside the
compiler. "Kernel wrapper" function contains body of SYCL kernel function,
receives OpenCL like parameters and additionally does some manipulation to
initialize captured lambda fields with this parameters. In some pseudo code
"kernel wrapper" looks like this:
// SYCL kernel is defined in SYCL headers
__kernel someSYCLKernel(lambda) {
lambda();
}
// Kernel wrapper
__kernel wrapper(global int* a) {
lambda; // Actually lambda declaration doesn't have a name in AST
// Let lambda has one captured field - accessor A. We need to init it with
global pointer from arguments:
lambda.A.__init(a);
// Body of SYCL kernel from SYCL headers:
{
lambda();
}
}
And actually kernel wrapper is presented in result module and passed to OpenCL
backend.
As I said, kernel wrapper is generated by the compiler inside the Sema and
OpenCLKernel attribute manually added to it, no matter which attribute was
added to "SYCL kernel" from SYCL headers.
So, while we are generating this wrapper I see only one profit to use OpenCL
kernel attribute in SYCL kernels - don't add new attribute to clang (but we
need to add `__kernel` keyword to SYCL).
I thought about idea - don't generate kernel wrapper but looks like it will not
work with OpenCL since we can't pass OpenCL `cl_mem` arguments inside any
structures (including accessors and lambdas) to the kernel.
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
