[SOLVED] Facing error: SYCL kernel cannot call a recursive function

Facing error: SYCL kernel cannot call a recursive function

I was running this code using SYCL and this error was coming up regarding recursion "error: SYCL kernel cannot call a recursive function" I am not sure what is causing this error. I used ordered set in the SYCL kernel and I am facing this issue not with the unordered set. I have mentioned the part of code below:

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]);//1st error source
            for (int i : a1[it]){ 
                i;
            }
        });

Error:

   
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

Solution

SYCL does not currently support recursive functions inside "device" code, i.e. a SYCL kernel. See the SYCL 2020 spec for confirmation.

This blog post has a section that talks about how to adapt code to avoid recursion using an example that might help under "Diffuse shading in the reference implementation - Recursive approach"