Thrust transform failure

[sdalton1]

The following code runs correctly on using host_vectors but fails at runtime using cuda device_vectors.

#include <thrust/device_vector.h>
#include <thrust/functional.h>
#include <thrust/sequence.h>
#include <thrust/transform.h>
#include <thrust/tuple.h>
#include <thrust/iterator/zip_iterator.h>

template<typename IndexType>
struct comp : public thrust::unary_function<thrust::tuple<IndexType,IndexType>,bool>
{
    template<typename Tuple>
    bool operator()(const Tuple& t)
    {   
        return thrust::get<0>(t) == thrust::get<1>(t);
    }   
};

int main(int argc, char*argv[])
{
    typedef thrust::device_vector<int> Vector;

    Vector x(10);
    Vector y(10);
    Vector z(10);

    thrust::sequence(x.begin(), x.end());
    thrust::sequence(y.begin(), y.end());

    thrust::transform(thrust::make_zip_iterator(thrust::make_tuple(x.begin(), y.begin())),
                      thrust::make_zip_iterator(thrust::make_tuple(x.end(), y.end())),
                      z.begin(),
                      comp<int>());

    return 0;
}

GDB backtrace

terminate called after throwing an instance of 'thrust::system::system_error'
  what():  bulk_kernel_by_value: an illegal memory access was encountered

Program received signal SIGABRT, Aborted.
0x00007ffff72d4cc9 in __GI_raise (sig=sig@entry=6) at ../nptl/sysdeps/unix/sysv/linux/raise.c:56
56      ../nptl/sysdeps/unix/sysv/linux/raise.c: No such file or directory.
(gdb) bt
#0  0x00007ffff72d4cc9 in __GI_raise (sig=sig@entry=6) at ../nptl/sysdeps/unix/sysv/linux/raise.c:56
#1  0x00007ffff72d80d8 in __GI_abort () at abort.c:89
#2  0x00007ffff78d96b5 in __gnu_cxx::__verbose_terminate_handler() () from /usr/lib/x86_64-linux-gnu/libstdc++.so.6
#3  0x00007ffff78d7836 in ?? () from /usr/lib/x86_64-linux-gnu/libstdc++.so.6
#4  0x00007ffff78d7863 in std::terminate() () from /usr/lib/x86_64-linux-gnu/libstdc++.so.6
#5  0x00007ffff78d7aa2 in __cxa_throw () from /usr/lib/x86_64-linux-gnu/libstdc++.so.6
#6  0x0000000000403b68 in thrust::system::cuda::detail::bulk_::detail::throw_on_error (e=cudaErrorIllegalAddress, message=0x40fdab "bulk_kernel_by_value") at /usr/local/cuda/include/thrust/system/cuda/detail/bulk/detail/throw_on_error.hpp:39
#7  0x0000000000404538 in thrust::system::cuda::detail::bulk_::detail::synchronize (message=0x40fdab "bulk_kernel_by_value") at /usr/local/cuda/include/thrust/system/cuda/detail/bulk/detail/synchronize.hpp:36
#8  0x0000000000404552 in thrust::system::cuda::detail::bulk_::detail::synchronize_if_enabled (message=0x40fdab "bulk_kernel_by_value") at /usr/local/cuda/include/thrust/system/cuda/detail/bulk/detail/synchronize.hpp:50
#9  0x0000000000408c5e in thrust::system::cuda::detail::bulk_::detail::cuda_launcher_base<0u, thrust::system::cuda::detail::bulk_::parallel_group<thrust::system::cuda::detail::bulk_::concurrent_group<thrust::system::cuda::detail::bulk_::agent<1ul>, 0ul>, 0ul>, thrust::system::cuda::detail::bulk_::detail::closure<thrust::system::cuda::detail::for_each_n_detail::for_each_kernel, thrust::tuple<thrust::system::cuda::detail::bulk_::detail::cursor<0u>, thrust::zip_iterator<thrust::tuple<thrust::zip_iterator<thrust::tuple<thrust::detail::normal_iterator<thrust::device_ptr<int> >, thrust::detail::normal_iterator<thrust::device_ptr<int> >, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type> >, thrust::detail::normal_iterator<thrust::device_ptr<int> >, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type> >, thrust::detail::wrapped_function<thrust::detail::unary_transform_functor<comp<int> >, void>, unsigned int, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type> > >::launch (this=0x7fffffffd820, num_blocks=1, block_size=1024, num_dynamic_smem_bytes=24544, stream=0x0, task=...) at /usr/local/cuda/include/thrust/system/cuda/detail/bulk/detail/cuda_launcher/cuda_launcher.hpp:74
#10 0x0000000000408217 in thrust::system::cuda::detail::bulk_::detail::cuda_launcher<thrust::system::cuda::detail::bulk_::parallel_group<thrust::system::cuda::detail::bulk_::concurrent_group<thrust::system::cuda::detail::bulk_::agent<1ul>, 0ul>, 0ul>, thrust::system::cuda::detail::bulk_::detail::closure<thrust::system::cuda::detail::for_each_n_detail::for_each_kernel, thrust::tuple<thrust::system::cuda::detail::bulk_::detail::cursor<0u>, thrust::zip_iterator<thrust::tuple<thrust::zip_iterator<thrust::tuple<thrust::detail::normal_iterator<thrust::device_ptr<int> >, thrust::detail::normal_iterator<thrust::device_ptr<int> >, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type> >, thrust::detail::normal_iterator<thrust::device_ptr<int> >, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type> >, thrust::detail::wrapped_function<thrust::detail::unary_transform_functor<comp<int> >, void>, unsigned int, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type> > >::launch (this=0x7fffffffd820, request=..., c=..., stream=0x0) at /usr/local/cuda/include/thrust/system/cuda/detail/bulk/detail/cuda_launcher/cuda_launcher.hpp:278
#11 0x0000000000407896 in thrust::system::cuda::detail::bulk_::detail::async_in_stream<thrust::system::cuda::detail::bulk_::parallel_group<thrust::system::cuda::detail::bulk_::concurrent_group<thrust::system::cuda::detail::bulk_::agent<1ul>, 0ul>, 0ul>, thrust::system::cuda::detail::bulk_::detail::closure<thrust::system::cuda::detail::for_each_n_detail::for_each_kernel, thrust::tuple<thrust::system::cuda::detail::bulk_::detail::cursor<0u>, thrust::zip_iterator<thrust::tuple<thrust::zip_iterator<thrust::tuple<thrust::detail::normal_iterator<thrust::device_ptr<int> >, thrust::detail::normal_iterator<thrust::device_ptr<int> >, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type> >, thrust::detail::normal_iterator<thrust::device_ptr<int> >, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type> >, thrust::detail::wrapped_function<thrust::detail::unary_transform_functor<comp<int> >, void>, unsigned int, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type> > > (g=..., c=..., s=0x0, before_event=0x0) at /usr/local/cuda/include/thrust/system/cuda/detail/bulk/detail/async.inl:46
#12 0x000000000040714d in thrust::system::cuda::detail::bulk_::detail::async<thrust::system::cuda::detail::bulk_::parallel_group<thrust::system::cuda::detail::bulk_::concurrent_group<thrust::system::cuda::detail::bulk_::agent<1ul>, 0ul>, 0ul>, thrust::system::cuda::detail::bulk_::detail::closure<thrust::system::cuda::detail::for_each_n_detail::for_each_kernel, thrust::tuple<thrust::system::cuda::detail::bulk_::detail::cursor<0u>, thrust::zip_iterator<thrust::tuple<thrust::zip_iterator<thrust::tuple<thrust::detail::normal_iterator<thrust::device_ptr<int> >, thrust::detail::normal_iterator<thrust::device_ptr<int> >, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type> >, thrust::detail::normal_iterator<thrust::device_ptr<int> >, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type> >, thrust::detail::wrapped_function<thrust::detail::unary_transform_functor<comp<int> >, void>, unsigned int, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type> > > (
    g=..., c=...) at /usr/local/cuda/include/thrust/system/cuda/detail/bulk/detail/async.inl:88
#13 0x0000000000406bcf in thrust::system::cuda::detail::bulk_::async<thrust::system::cuda::detail::bulk_::parallel_group<thrust::system::cuda::detail::bulk_::concurrent_group<thrust::system::cuda::detail::bulk_::agent<1ul>, 0ul>, 0ul>, thrust::system::cuda::detail::for_each_n_detail::for_each_kernel, thrust::system::cuda::detail::bulk_::detail::cursor<0u>, thrust::zip_iterator<thrust::tuple<thrust::zip_iterator<thrust::tuple<thrust::detail::normal_iterator<thrust::device_ptr<int> >, thrust::detail::normal_iterator<thrust::device_ptr<int> >, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type> >, thrust::detail::normal_iterator<thrust::device_ptr<int> >, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type> >, thrust::detail::wrapped_function<thrust::detail::unary_transform_functor<comp<int> >, void>, unsigned int> (g=..., f=..., arg1=..., arg2=..., arg3=..., arg4=10) at /usr/local/cuda/include/thrust/system/cuda/detail/bulk/detail/async.inl:141
#14 0x000000000040650e in thrust::zip_iterator<thrust::tuple<thrust::zip_iterator<thrust::tuple<thrust::detail::normal_iterator<thrust::device_ptr<int> >, thrust::detail::normal_iterator<thrust::device_ptr<int> >, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type> >, thrust::detail::normal_iterator<thrust::device_ptr<int> >, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type> > thrust::system::cuda::detail::for_each_n<thrust::system::cuda::detail::tag, thrust::zip_iterator<thrust::tuple<thrust::zip_iterator<thrust::tuple<thrust::detail::normal_iterator<thrust::device_ptr<int> >, thrust::detail::normal_iterator<thrust::device_ptr<int> >, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type> >, thrust::detail::normal_iterator<thrust::device_ptr<int> >, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type> >, long, thrust::detail::unary_transform_functor<comp<int> > >(thrust::system::cuda::detail::execution_policy<thrust::system::cuda::detail::tag>&, thrust::zip_iterator<thrust::tuple<thrust::zip_iterator<thrust::tuple<thrust::detail::normal_iterator<thrust::device_ptr<int> >, thrust::detail::normal_iterator<thrust::device_ptr<int> >, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type> >, thrust::detail::normal_iterator<thrust::device_ptr<int> >, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type> >, long, thrust::detail::unary_transform_functor<comp<int> >)::workaround::parallel_path(thrust::system::cuda::detail::execution_policy<thrust::system::cuda::detail::tag>&, thrust::zip_iterator<thrust::tuple<thrust::zip_iterator<thrust::tuple<thrust::detail::normal_iterator<thrust::device_ptr<int> >, thrust::detail::normal_iterator<thrust::device_ptr<int> >, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type> >, thrust::detail::normal_iterator<thrust::device_ptr<int> >, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type> >, long, thrust::detail::unary_transform_functor<comp<int> >) (exec=..., first=..., n=10, f=...) at /usr/local/cuda/include/thrust/system/cuda/detail/for_each.inl:143
#15 0x0000000000406599 in thrust::system::cuda::detail::for_each_n<thrust::system::cuda::detail::tag, thrust::zip_iterator<thrust::tuple<thrust::zip_iterator<thrust::tuple<thrust::detail::normal_iterator<thrust::device_ptr<int> >, thrust::detail::normal_iterator<thrust::device_ptr<int> >, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type> >, thrust::detail::normal_iterator<thrust::device_ptr<int> >, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type> >, long, thrust::detail::unary_transform_functor<comp<int> > > (exec=..., 
    first=..., n=10, f=...) at /usr/local/cuda/include/thrust/system/cuda/detail/for_each.inl:157
#16 0x0000000000405e98 in thrust::system::cuda::detail::for_each<thrust::system::cuda::detail::tag, thrust::zip_iterator<thrust::tuple<thrust::zip_iterator<thrust::tuple<thrust::detail::normal_iterator<thrust::device_ptr<int> >, thrust::detail::normal_iterator<thrust::device_ptr<int> >, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type> >, thrust::detail::normal_iterator<thrust::device_ptr<int> >, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type> >, thrust::detail::unary_transform_functor<comp<int> > > (exec=..., 
    first=..., last=..., f=...) at /usr/local/cuda/include/thrust/system/cuda/detail/for_each.inl:173

[schiller-manuel]

I can confirm this failure with CUDA 6, 6.5. and 7 and the latest thrust version from github.

[jaredhoberock]

Your functor is missing something, Steven :-)

[sdalton1]

Doh! host device decorators. Thanks Jared!

[schiller-manuel]

ah, didn't see that either. but why did it even compile?