将lambda函数传递给区域的SYCL并行

Question

SYCL内核可以提交给queue，如下所示-

Queue.submit([=](handler& cgh){
        cgh.parallel_for<class test>(_range, [=](nd_item<2> iter{
      Write your function (F) here. 
});
    });

现在，上面提到的函数F被接受为lambda，我想这样做如下-

void deviceManager::dispatchFunction(const T CGF, nd_range<2> _range){
    this->Queue.submit([=](handler& cgh){
        cgh.parallel_for<class test>(_range, CGF);
    });

其中CGF是作为参数传递的lamda函数。我是这样称呼的-

    int y = 0;
    int z = 10;
    int x = -2394801;

    deviceManager manager;   
    manager.dispatchFunction([=]() -> void {
        printf("Printf from lambda  | %d = \n", x);
        printf("Y = %d | Z = %d \n", y, y + z);
    }, nd_range<2>(range<2>(4, 4), range<2>(2, 2)));

但是，我得到了以下错误-

In file included from testSYCL.cpp:4:
In file included from /opt/intel/oneapi/compiler/2022.1.0/linux/bin-llvm/../include/sycl/CL/sycl.hpp:16:
In file included from /opt/intel/oneapi/compiler/2022.1.0/linux/bin-llvm/../include/sycl/CL/sycl/backend.hpp:18:
In file included from /opt/intel/oneapi/compiler/2022.1.0/linux/bin-llvm/../include/sycl/CL/sycl/detail/backend_traits_opencl.hpp:25:
In file included from /opt/intel/oneapi/compiler/2022.1.0/linux/bin-llvm/../include/sycl/CL/sycl/queue.hpp:20:
/opt/intel/oneapi/compiler/2022.1.0/linux/bin-llvm/../include/sycl/CL/sycl/handler.hpp:1112:5: error: no matching function for call to object of type 'const (lambda at testSYCL.cpp:141:30)'
    KernelFunc(detail::Builder::getElement(detail::declptr<ElementType>()));
    ^~~~~~~~~~
/opt/intel/oneapi/compiler/2022.1.0/linux/bin-llvm/../include/sycl/CL/sycl/handler.hpp:1232:5: note: in instantiation of function template specialization 'sycl::handler::kernel_parallel_for<test, sycl::nd_item<2>, (lambda at testSYCL.cpp:141:30)>' requested here
    kernel_parallel_for<KernelName, ElementType>(KernelFunc);
    ^
/opt/intel/oneapi/compiler/2022.1.0/linux/bin-llvm/../include/sycl/CL/sycl/handler.hpp:1568:5: note: in instantiation of function template specialization 'sycl::handler::kernel_parallel_for_wrapper<test, sycl::nd_item<2>, (lambda at testSYCL.cpp:141:30)>' requested here
    kernel_parallel_for_wrapper<NameT, TransformedArgType>(KernelFunc);
    ^
testSYCL.cpp:124:13: note: in instantiation of function template specialization 'sycl::handler::parallel_for<test, (lambda at testSYCL.cpp:141:30), 2>' requested here
        cgh.parallel_for<class test>(_range, CGF);
            ^
testSYCL.cpp:141:13: note: in instantiation of function template specialization 'deviceManager::dispatchFunction<(lambda at testSYCL.cpp:141:30)>' requested here
    manager.dispatchFunction([=]() -> void {
            ^
testSYCL.cpp:141:30: note: candidate function not viable: requires 0 arguments, but 1 was provided
    manager.dispatchFunction([=]() -> void {
                             ^
1 error generated.

怎样才能达到预期的效果？

提前谢谢

Answer 1

传递给sycl::handler::parallel_for()的SYCL内核函数应该至少有一个index space class参数。因此，正如注释中提到的@Artyer和@Slava，lambda函数至少应该有一个参数，即使它没有被使用：

    manager.dispatchFunction([=](cl::sycl::nd_item<2>) -> void {
        printf("Printf from lambda  | %d = \n", x);
        printf("Y = %d | Z = %d \n", y, y + z);
    }, nd_range<2>(range<2>(4, 4), range<2>(2, 2)));

接下来，SYCL不支持设备代码中的printf()函数调用。有必要使用sycl::stream代替它(请参阅C.10.4.printf函数部分)。下面是一个简短的例子，它基于问题中的代码：

class DeviceManager {
public:
    DeviceManager(const sycl::device_selector& d_selector)
        : m_queue(d_selector)
    {}

    template<typename Func, int N>
    void dispatchFunction(Func func, const cl::sycl::nd_range<N>& range) {
        m_queue.submit([=](cl::sycl::handler& cgh) {
            cl::sycl::stream out(1024, 256, cgh);
            cgh.parallel_for(range, [=](cl::sycl::nd_item<N> iter) {
                std::invoke(func, iter, out);
            });
        });
    }

    void wait() {
        m_queue.wait();
    }

private:
    cl::sycl::queue m_queue;
};

int main() {
    DeviceManager deviceManager((cl::sycl::default_selector()));

    int y = 0;
    int z = 10;
    int x = -2394801;

    deviceManager.dispatchFunction([=](cl::sycl::nd_item<2>, cl::sycl::stream sout) {
        sout << "Printf from lambda  | " << x << " = \n";
        sout << "Y = " << y << " | Z = " << y + z << "\n";
    }, cl::sycl::nd_range<2>(cl::sycl::range<2>(4, 4), cl::sycl::range<2>(2, 2)));
    deviceManager.wait();
}