SYCL内核使用ordered_set触发递归函数调用报错求助

问题根源 #

SYCL内核代码禁止调用递归函数，这是因为GPU等加速器设备的栈空间远小于CPU，递归极易引发栈溢出；同时SYCL设备编译器对递归函数的编译优化支持有限，因此会直接拦截这类调用。

你遇到的问题本质是：

• std::ordered_set底层依赖红黑树实现，红黑树的核心操作（如复制_M_copy、删除_M_erase）采用了递归写法。当你在SYCL内核中访问ordered_set（比如调用foo(a1[it])或遍历集合）时，会间接触发这些递归函数，触发编译报错。
• std::unordered_set底层是哈希表结构，所有核心操作都是非递归实现，因此不会触发SYCL的递归限制。

解决方案 #

• 改用非递归容器：如果业务逻辑不需要有序集合，直接使用std::unordered_set是最简便的方案。
• 将集合操作移至主机端：若必须使用有序集合，把对ordered_set的处理（比如foo函数内的逻辑、遍历操作）放在SYCL内核之外的主机代码中，仅将需要设备计算的数据传入内核。
• 使用SYCL兼容的有序容器：部分SYCL实现（如Intel oneAPI DPC++）提供了设备优化的容器库，比如oneapi::tbb::ordered_set，这类容器专为设备端设计，避免了递归实现，可直接在内核中使用。
• 手动实现非递归红黑树操作：如果需要自定义有序集合逻辑，可以重写红黑树的核心方法，用循环替代递归，确保代码能通过SYCL设备编译器的检查。

问题代码片段 #

sycl::buffer b1(sets);
q.submit([&](sycl::handler& h) {
    sycl::accessor a1(b1,h);
    h.parallel_for(sycl::range<1>(sets.size()), [=](sycl::item<1> it) {
        foo(a1[it]);// 错误触发点
        for (int i : a1[it]){
            i;
        }
    });
});

错误信息 #

In file included from /usr/lib/gcc/x86_64-linux-gnu/9/../../../../include/c++/9/set:60:
/usr/lib/gcc/x86_64-linux-gnu/9/../../../../include/c++/9/bits/stl_tree.h:1886:26: error: SYCL kernel cannot call a recursive function
              __top->_M_right = _M_copy(_S_right(__x), __top, __node_gen);
                                ^
/usr/lib/gcc/x86_64-linux-gnu/9/../../../../include/c++/9/bits/stl_tree.h:1875:7: note: function implemented using recursion declared here
      typename _Rb_tree<_Key, _Val, _KoV, _Compare, _Alloc>::_Link_type
      ^
/usr/lib/gcc/x86_64-linux-gnu/9/../../../../include/c++/9/bits/stl_tree.h:1896:21: error: SYCL kernel cannot call a recursive function
                  __y->_M_right = _M_copy(_S_right(__x), __y, __node_gen);
                                  ^
/usr/lib/gcc/x86_64-linux-gnu/9/../../../../include/c++/9/bits/stl_tree.h:1875:7: note: function implemented using recursion declared here
      typename _Rb_tree<_Key, _Val, _KoV, _Compare, _Alloc>::_Link_type
      ^
/usr/lib/gcc/x86_64-linux-gnu/9/../../../../include/c++/9/bits/stl_tree.h:1903:6: error: SYCL kernel cannot call a recursive function
            _M_erase(__top);
            ^
/usr/lib/gcc/x86_64-linux-gnu/9/../../../../include/c++/9/bits/stl_tree.h:1911:5: note: function implemented using recursion declared here
    void
    ^
/usr/lib/gcc/x86_64-linux-gnu/9/../../../../include/c++/9/bits/stl_tree.h:1000:9: error: SYCL kernel cannot call a recursive function
      { _M_erase(_M_begin()); }
        ^
/usr/lib/gcc/x86_64-linux-gnu/9/../../../../include/c++/9/bits/stl_tree.h:1911:5: note: function implemented using recursion declared here
    void
    ^
/usr/lib/gcc/x86_64-linux-gnu/9/../../../../include/c++/9/bits/stl_tree.h:1918:4: error: SYCL kernel cannot call a recursive function
          _M_erase(_S_right(__x));
          ^
/usr/lib/gcc/x86_64-linux-gnu/9/../../../../include/c++/9/bits/stl_tree.h:1911:5: note: function implemented using recursion declared here
    void
    ^
/usr/lib/gcc/x86_64-linux-gnu/9/../../../../include/c++/9/bits/stl_tree.h:906:24: error: SYCL kernel cannot call a recursive function
          _Link_type __root = _M_copy(__x._M_begin(), _M_end(), __gen);
                              ^
/usr/lib/gcc/x86_64-linux-gnu/9/../../../../include/c++/9/bits/stl_tree.h:1875:7: note: function implemented using recursion declared here
      typename _Rb_tree<_Key, _Val, _KoV, _Compare, _Alloc>::_Link_type
      ^
/usr/lib/gcc/x86_64-linux-gnu/9/../../../../include/c++/9/bits/stl_tree.h:366:12: error: SYCL kernel cannot call an undefined function without SYCL_EXTERNAL attribute
        _M_node = _Rb_tree_increment(_M_node);
                  ^
/usr/lib/gcc/x86_64-linux-gnu/9/../../../../include/c++/9/bits/stl_tree.h:247:3: note: '_Rb_tree_increment' declared here
  _Rb_tree_increment(const _Rb_tree_node_base* __x) throw ();
  ^
/usr/lib/gcc/x86_64-linux-gnu/9/../../../../include/c++/9/bits/stl_tree.h:364:7: note: called by 'operator++'
      operator++() _GLIBCXX_NOEXCEPT
      ^
7 errors generated.
make: *** [<builtin>: test_set] Error 1

内容的提问来源于stack exchange，提问作者Pksingh