TestZipIteratorReduceByKey is broken with sm_11

[GoogleCodeExporter]

jhoberock@jhoberock-dt:~/dev/hg$ hg clone https://thrust.googlecode.com/hg 
thrust-zip-iterator-reduce-by-key-bug
requesting all changes
adding changesets
adding manifests
adding file changes
added 1879 changesets with 10994 changes to 2561 files (+7 heads)
updating to branch default
813 files updated, 0 files merged, 0 files removed, 0 files unresolved
jhoberock@jhoberock-dt:~/dev/hg$ cd thrust-zip-iterator-reduce-by-key-bug/
jhoberock@jhoberock-dt:~/dev/hg/thrust-zip-iterator-reduce-by-key-bug$ cd 
testing/
jhoberock@jhoberock-dt:~/dev/hg/thrust-zip-iterator-reduce-by-key-bug/testing$ 
scons arch=sm_11 tests=ZipIteratorReduceByKey -j4
scons: Reading SConscript files ...
scons: done reading SConscript files.
scons: Building targets ...
nvcc -o testframework.o -c -arch=sm_11 -Xcompiler 
-DTHRUST_DEVICE_SYSTEM=THRUST_DEVICE_SYSTEM_CUDA -Xcompiler 
-DTHRUST_HOST_SYSTEM=THRUST_HOST_SYSTEM_CPP -Xcompiler -O2 -Xcompiler -Wall -I 
/home/jhoberock/dev/hg/thrust-zip-iterator-reduce-by-key-bug -I 
/usr/local/cuda/include -I 
/home/jhoberock/dev/hg/thrust-zip-iterator-reduce-by-key-bug/testing 
testframework.cu
nvcc -o zip_iterator_reduce_by_key.o -c -arch=sm_11 -Xcompiler 
-DTHRUST_DEVICE_SYSTEM=THRUST_DEVICE_SYSTEM_CUDA -Xcompiler 
-DTHRUST_HOST_SYSTEM=THRUST_HOST_SYSTEM_CPP -Xcompiler -O2 -Xcompiler -Wall -I 
...
gcc -o tester testframework.o zip_iterator_reduce_by_key.o 
-L/usr/local/cuda/lib64 -lstdc++ -lcudart
scons: done building targets.
jhoberock@jhoberock-dt:~/dev/hg/thrust-zip-iterator-reduce-by-key-bug/testing$ 
./tester 
Running 1 unit tests.
F
================================================================
FAILURE: TestZipIteratorReduceByKey
[zip_iterator_reduce_by_key.cu:87] Sequences are not equal [type='long']
--------------------------------
  [536] 6042755086834252800  5073566479746842624
  [788] -704063934970552320  -6253648293480741888
--------------------------------
Sequences differ at 2 of 1565 positions

================================================================
Totals: 1 failures, 0 known failures, 0 errors, and 0 passes.
Time:  0.0166667 minutes

jhoberock@jhoberock-dt:~/dev/hg/thrust-zip-iterator-reduce-by-key-bug/testing$ 
cat /proc/driver/nvidia/gpus/0/information 
Model:       GeForce 8800 GT
IRQ:         16
Video BIOS:      62.92.12.00.04
Card Type:   PCI-E
DMA Size:    40 bits
DMA Mask:    0xffffffffff
Bus Location:    0000:01.00.0

Original issue reported on code.google.com by jaredhoberock on 1 Mar 2012 at 12:34

[GoogleCodeExporter]

Here's a reproducer for the smallest N i could find:

#include <unittest/unittest.h>
#include <thrust/iterator/zip_iterator.h>
#include <thrust/reduce.h>
#include <thrust/host_vector.h>
#include <thrust/device_vector.h>

template<typename Tuple>
struct TuplePlus
{
  __host__ __device__
  Tuple operator()(Tuple x, Tuple y) const
  {
    using namespace thrust;
    return make_tuple(get<0>(x) + get<0>(y),
                      get<1>(x) + get<1>(y));
  }
}; // end TuplePlus

int main()
{
  size_t n = 1505;

  using namespace thrust;

  typedef long T;

  host_vector<T> h_data0 = unittest::random_integers<bool>(n);
  host_vector<T> h_data1 = unittest::random_integers<T>(n);
  host_vector<T> h_data2 = unittest::random_integers<T>(n);
  host_vector<T> h_data3(n,0);
  host_vector<T> h_data4(n,0);
  host_vector<T> h_data5(n,0);

  device_vector<T> d_data0 = h_data0;
  device_vector<T> d_data1 = h_data1;
  device_vector<T> d_data2 = h_data2;
  device_vector<T> d_data3(n,0);
  device_vector<T> d_data4(n,0);
  device_vector<T> d_data5(n,0);

  typedef tuple<T,T> Tuple;

  // run on host
  reduce_by_key
      ( make_zip_iterator(make_tuple(h_data0.begin(), h_data0.begin())),
        make_zip_iterator(make_tuple(h_data0.end(),   h_data0.end())),
        make_zip_iterator(make_tuple(h_data1.begin(), h_data2.begin())),
        make_zip_iterator(make_tuple(h_data3.begin(), h_data3.begin())),
        make_zip_iterator(make_tuple(h_data4.begin(), h_data5.begin())),
        equal_to<Tuple>(),
        TuplePlus<Tuple>());

  // run on device
  reduce_by_key
      ( make_zip_iterator(make_tuple(d_data0.begin(), d_data0.begin())),
        make_zip_iterator(make_tuple(d_data0.end(),   d_data0.end())),
        make_zip_iterator(make_tuple(d_data1.begin(), d_data2.begin())),
        make_zip_iterator(make_tuple(d_data3.begin(), d_data3.begin())),
        make_zip_iterator(make_tuple(d_data4.begin(), d_data5.begin())),
        equal_to<Tuple>(),
        TuplePlus<Tuple>());

  ASSERT_EQUAL(h_data3, d_data3);
  ASSERT_EQUAL(h_data4, d_data4);
  ASSERT_EQUAL(h_data5, d_data5);

  return 0;
}

Original comment by jaredhoberock on 1 Mar 2012 at 9:59

[GoogleCodeExporter]

This issue was closed by revision 257c6c3a96a6.

Original comment by wnbell on 2 Mar 2012 at 10:31

• Changed state: Fixed