Cannot access copied lambda vectors
Closed this issue · 2 comments
domantasjurkus commented
I am having problems with the following bit of code, namely the two locations where in is accessed within the algorithm call:
void f(std::vector<float> &out, std::vector<float> in, uint2 outSize) {
std::vector<uint> pixels(outSize.x*outSize.y);
std::iota(pixels.begin(), pixels.end(), 0);
std::experimental::parallel::transform(half_sample_par, pixels.begin(), pixels.end(), out.begin(), [=](uint pos) {
uint x = pos % outSize.x;
uint y = pos / outSize.x;
const float center = in[centerPixel.x + centerPixel.y*inSize.x];
std::vector<int2> pairs = generate_int_pairs(-r+1, r, -r+1, r);
std::for_each(pairs.begin(), pairs.end(), [&](int2 p) {
...
float current = in[cur.x + cur.y*inSize.x];
...
});
return t/sum;
});
}
I get the error below during compile time. I understand that in is essentially a vector of pointers and that the reference copied by the lambda cannot be used to access memory. Is there any easy workaround for this? For all the algorithm is concerned, in is a static variable used only for reading in values.
Using two input iterators for std::transform is not an option as in is indexed at different locations in the function. I am fairly sure zipping the in values along with pixels will not help much.
[ 71%] Building ComputeCpp integration header file /home/dom/gu/Project/slambench/build/kfusion/track.cpp.sycl
/home/dom/gu/Project/slambench/kfusion/src/stl/track.cpp:33:30: error: [Computecpp:CC0011]: cannot capture object in of type 'std::vector<float>' in a SYCL kernel, because it contains a field with a
pointer type
const float center = in[centerPixel.x + centerPixel.y*inSize.x];
^
/home/dom/computecpp-ce-0.5.1/include/SYCL/apis.h:1299:5: note: in instantiation of function template specialization 'cl::sycl::kernelgen_parallel_for_nd<(anonymous namespace)::half_sample, (lambda
at /home/dom/sycl_parallel_stl/include/sycl/algorithm/transform.hpp:71:14)>' requested here
kernelgen_parallel_for_nd<
^
/home/dom/sycl_parallel_stl/include/sycl/algorithm/transform.hpp:70:9: note: in instantiation of function template specialization 'cl::sycl::handler::parallel_for<(anonymous namespace)::half_sample,
(lambda at /home/dom/sycl_parallel_stl/include/sycl/algorithm/transform.hpp:71:14), 1>' requested here
h.parallel_for<typename ExecutionPolicy::kernelName>(
^
/home/dom/sycl_parallel_stl/include/sycl/execution_policy:202:18: note: in instantiation of function template specialization 'sycl::impl::transform<sycl::sycl_execution_policy<(anonymous
namespace)::half_sample>, __gnu_cxx::__normal_iterator<unsigned int *, std::vector<unsigned int, std::allocator<unsigned int> > >, __gnu_cxx::__normal_iterator<float *, std::vector<float,
std::allocator<float> > >, (lambda at /home/dom/gu/Project/slambench/kfusion/src/stl/track.cpp:23:104)>' requested here
return impl::transform(named_sep, b, e, out_b, op);
^
/home/dom/sycl_parallel_stl/include/experimental/algorithm:132:14: note: in instantiation of function template specialization 'sycl::sycl_execution_policy<(anonymous
namespace)::half_sample>::transform<__gnu_cxx::__normal_iterator<unsigned int *, std::vector<unsigned int, std::allocator<unsigned int> > >, __gnu_cxx::__normal_iterator<float *,
std::vector<float, std::allocator<float> > >, (lambda at /home/dom/gu/Project/slambench/kfusion/src/stl/track.cpp:23:104)>' requested here
return sep.transform(b, e, out, op);
^
/home/dom/gu/Project/slambench/kfusion/src/stl/track.cpp:23:34: note: in instantiation of function template specialization
'std::experimental::parallel::transform<sycl::sycl_execution_policy<(anonymous namespace)::half_sample> &, __gnu_cxx::__normal_iterator<unsigned int *, std::vector<unsigned int,
std::allocator<unsigned int> > >, __gnu_cxx::__normal_iterator<float *, std::vector<float, std::allocator<float> > >, (lambda at
/home/dom/gu/Project/slambench/kfusion/src/stl/track.cpp:23:104)>' requested here
std::experimental::parallel::transform(half_sample_par, pixels.begin(), pixels.end(), out.begin(), [=](uint pos) {
^
/home/dom/gu/Project/slambench/kfusion/src/stl/track.cpp:12:9: note: captured object defined here
std::vector<float> in,
^
/usr/bin/../lib/gcc/x86_64-linux-gnu/7.2.0/../../../../include/c++/7.2.0/bits/stl_vector.h:216:20: note: std::vector derives from std::_Vector_base
class vector : protected _Vector_base<_Tp, _Alloc>
^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/usr/bin/../lib/gcc/x86_64-linux-gnu/7.2.0/../../../../include/c++/7.2.0/bits/stl_vector.h:166:20: note: std::_Vector_base contains field std::_Vector_base<float, std::allocator<float> >::_M_impl of
type 'struct std::_Vector_base<float, class std::allocator<float> >::_Vector_impl'
_Vector_impl _M_impl;
~~~~~~~~~~~~~^~~~~~~
/usr/bin/../lib/gcc/x86_64-linux-gnu/7.2.0/../../../../include/c++/7.2.0/bits/stl_vector.h:84:2: note: offending field std::_Vector_base<float, std::allocator<float> >::_Vector_impl::_M_start of
type 'pointer' defined here
pointer _M_start;
^
In file included from /home/dom/gu/Project/slambench/kfusion/src/stl/track.cpp:1:
In file included from /home/dom/gu/Project/slambench/kfusion/./include/kernels_stl.h:5:
In file included from /home/dom/gu/Project/slambench/kfusion/./include/commons.h:35:
In file included from /home/dom/gu/Project/slambench/kfusion/./include/default_parameters.h:15:
In file included from /usr/bin/../lib/gcc/x86_64-linux-gnu/7.2.0/../../../../include/c++/7.2.0/vector:69:
/usr/bin/../lib/gcc/x86_64-linux-gnu/7.2.0/../../../../include/c++/7.2.0/bits/vector.tcc:444:4: error: [Computecpp:CC0006]: Using exceptions is not allowed within SYCL device code
__throw_exception_again;
^
/usr/bin/../lib/gcc/x86_64-linux-gnu/7.2.0/../../../../include/c++/7.2.0/bits/exception_defines.h:42:34: note: expanded from macro '__throw_exception_again'
# define __throw_exception_again throw
^
In file included from /home/dom/gu/Project/slambench/kfusion/src/stl/track.cpp:1:
In file included from /home/dom/gu/Project/slambench/kfusion/./include/kernels_stl.h:5:
In file included from /home/dom/gu/Project/slambench/kfusion/./include/commons.h:35:
In file included from /home/dom/gu/Project/slambench/kfusion/./include/default_parameters.h:15:
In file included from /usr/bin/../lib/gcc/x86_64-linux-gnu/7.2.0/../../../../include/c++/7.2.0/vector:69:
/usr/bin/../lib/gcc/x86_64-linux-gnu/7.2.0/../../../../include/c++/7.2.0/bits/vector.tcc:408:7: error: [Computecpp:CC0006]: Using exceptions is not allowed within SYCL device code
__try
^
/usr/bin/../lib/gcc/x86_64-linux-gnu/7.2.0/../../../../include/c++/7.2.0/bits/exception_defines.h:40:21: note: expanded from macro '__try'
# define __try try
^
3 errors generated.
kfusion/CMakeFiles/kfusion-stl_track.cpp_2_ih.dir/build.make:61: recipe for target 'kfusion/track.cpp.sycl' failed
make[3]: *** [kfusion/track.cpp.sycl] Error 1
make[3]: Leaving directory '/home/dom/gu/Project/slambench/build'
CMakeFiles/Makefile2:555: recipe for target 'kfusion/CMakeFiles/kfusion-stl_track.cpp_2_ih.dir/all' failed
make[2]: *** [kfusion/CMakeFiles/kfusion-stl_track.cpp_2_ih.dir/all] Error 2
make[2]: Leaving directory '/home/dom/gu/Project/slambench/build'
Makefile:83: recipe for target 'all' failed
make[1]: *** [all] Error 2
make[1]: Leaving directory '/home/dom/gu/Project/slambench/build'
Makefile:14: recipe for target 'build' failed
make: *** [build] Error 2
tremmelg commented
I'm afraid that is a limitation of SYCL. You can't capture a std::vector in the lambda that will be your kernel, because it has a pointer field. You should try to make in an argument to std::experimental::parallel::transform, such that the appropriate element gets passed to the lambda along with the pixel id.
tremmelg commented
Oh, sorry, just saw that using two iterators is not an option.