Before, building kernels with strings was available using sycl::program::build_with_source and sycl::program::compile_with_source, but this is deprecated.
If I have a function defined in an std::string:
"int foo(int a, int b){return a+b}"
How can I compile and link this to the application, so that it becomes callable inside a kernel:
q.submit([&](sycl::handler& h){
auto accessor_read = ...
auto accessor_write = ...
h.parallel_for(N_tasks, [=](sycl::id id){
accessor_write[id] = foo(accessor_read[2*id], accessor_read[2*id+1]);
});
});
If this is impractical/impossible, I’m up for alternative solutions.
So the functionality that you’re looking for was originally in SYCL 1.2.1, however, it was moved into the OpenCL backend specification in SYCL 2020, as it’s an OpenCL backend-specific feature. I think the function you’re looking for is described in SYCL™ 2020 Specification (revision 7). Though note this is an optional feature so support across SYCL implementations may vary. This also assumes that you’re targeting the OpenCL backend.
Other options could be to create the program/kernel object from source via the native API and then use backend interoperability (SYCL™ 2020 Specification (revision 7)) to construct a kernel_bundle in SYCL.
Another way, assuming there is a SYCL implementation supporting dynamic library, would be to run from the main executable the SYCL compiler on a file with a host function submitting the command group you show with also the definition of foo in the same file to generate a dynamic library.
Then load the dynamic library and call the generated function with the queue and buffers as arguments.
Not tested on my side.
If you get this to work, make an open-source library doing this.
Dynamic linking of device functions is apparently a work in progress problem according to Intel-DPCPP:intel.github.io/llvm-docs/design/SharedLibraries.html
I tried compiling shared libraries by invoking the dpcpp-compiler with std::system, but SYCL does not accept the function pointers retrieved by dlopen.
Initializing a sycl::kernel with an OpenCL-kernel works, but brings back the challenges with target specification (which was solved for SYCL using -fsycl-targets). By building the OpenCL-kernel beforehand, I guess I’m no longer able to transpile to nvptx?
Thank you for trying!
I am unsure by what you mean with “SYCL does not accept the function pointers retrieved by dlopen.” In that part it is just about pure C++ and you just get a pointer to a normal C++ function (containing some SYCL code inside, put this should not be important at this point) and call it. Perhaps Boost.DLL Chapter 12. Boost.DLL - 1.82.0 can help on this side.
But the problem will probably happen when you execute the function and the function tries to launch the kernel (I can imagine an inclusion header mismatch).
Otherwise the OpenCL path you suggest used to be supported by ComputeCpp.
Nowadays, the path I can think of would requires PoCL to expose the Nvidia devices as OpenCL SPIR-V devices and redirecting to the CUDA + PTX one. PoCL - Portable Computing Language | NVIDIA GPU support via CUDA backend
