Hello,
I’ve been looking to find a way to pass an arbitrary number of accessors to the kernel but I’ve been having trouble with it. The code compiles and executes properly with the Host device, however it throws an exception during runtime on CL devices.
The general idea is to receive a variable number of iterators through a template to then create a buffer and accessor for each. Finally, an operation will be applied with the received data and stored somewhere else.
The general structure of the code is as follows:
First the iterators are passed by means of a tuple to the function.
template <typename ... Input, typename Output, typename Operation>
void my_function(
std::tuple<Input...> in_iters,
Output out,
Operation && op) const
{
Then each buffer for each iterator is created and stored in an std::array.
std::array in_buffers = {std::apply([size](const auto&... inputs){
std::array collection{sycl::buffer<T,1>{inputs, inputs + size}...};
return collection;
}, in_iters)};
A similar approach is taken with the creation of the accessors.
std::array in_accs = {std::apply([&] (auto&... buffers) {
std::array accessors{buffers.template get_access<sycl::access::mode::read>(cgh)...};
return accessors;
},in_buffers)};
Finally we have the kernel code that calls the lambda function with all of the other accessors and stores the result.
cgh.template parallel_for<myKernel>(sycl::range<1>{size}, [=] (sycl::id<1> index) {
out_acc[index] = std::apply([&](const auto &...accessors){
return op(accessors[index]...);
}, in_accs););
When run on a device different than the host, the error CL_INVALID_ARG_SIZE is caught. The problem occurs when ‘in_accs’ is used. I supposed that it’s because the variable is not an accessor but an array. I’ve also tried to use an accessor array (sycl::accessor[.]) with no luck. I’ve seen that DPC++ can execute this code but I’m unsure if it’s because of an extension in the compiler. I was wondering if there’s another approach to this since access to all of the variables at the same time is required due to the way the lambda function is structured.
Thank you in advance