CUDA_EXCEPTION_10: zip_iterator<device_vector> -> transform_iterator -> device_vector<tuple>

[GoogleCodeExporter]

When I cuda-gdb thrust programs with device_vector<tuple> derived from 
zip_iterator<tuple> through transform_iterator<tuple> compiled with the -G 
option, the programs stop printing "Program received signal CUDA_EXCEPTION_10, 
Device Illegal Address." This error is replayed by the following example:

#include <thrust/device_vector.h>
#include <thrust/host_vector.h>
#include <thrust/tuple.h>
#include <thrust/iterator/transform_iterator.h>
#include <thrust/iterator/zip_iterator.h>

#define N 1

int
main()
{
  thrust::device_vector<float> iv(N, 1.0f);
  typedef thrust::device_vector<float>::iterator i_t;
  typedef thrust::zip_iterator<thrust::tuple<i_t> > z_t;
  z_t z(thrust::make_tuple(iv.begin()));
  typedef thrust::tuple<float> t_t;
  typedef thrust::identity<t_t> id_t;
  thrust::transform_iterator<id_t, z_t> t(z, id_t());
  thrust::device_vector<t_t> ov(t, t + N);
  thrust::host_vector<t_t> hv(ov);
  return 0;
}

I tested this bug with Thrust 1.5.1 and 1.6.0, CUDA 4.1, GCC 4.1.2 and 4.4.6 
(specified as an NVCC's backend with the -ccbin option,) CentOS 5.8 x86_64 and 
Tesla C2050 with driver 295.33.

Original issue reported on code.google.com by ryo-nish...@hpc.co.jp on 26 Mar 2012 at 3:53

[GoogleCodeExporter]

This is a known compiler bug

Original comment by jaredhoberock on 7 May 2012 at 9:59

• Changed state: Invalid

[GoogleCodeExporter]

This bug has seemed to be fixed in the CUDA 4.2 compiler. Thank you for your 
investigating.

Original comment by ryo-nish...@hpc.co.jp on 9 May 2012 at 2:06