Reduce functors behave unexpectedly when argument types differ from standard

[GoogleCodeExporter]

  template<class R, class T, class U>
  struct Plus
  {
        __host__ __device__
        R operator()(const T& t, const U& u){
                return ((R)t) + ((R)u);
        }
  };

  typedef unsigned char type;
  thrust::device_vector<type> vec(1000,1);
  Plus<int,int,type> p;
  int sum    = thrust::reduce(vec.begin(), vec.end(), (int) 0, p);

surprisingly for type=unsigned char, this does not yield 1000 but 232.
I expect the binary_op to be used as in 

  sum=binary_op(sum,nextval),

which I confirmed in thrust source. Somewhere along the way there is a type 
issue which I coulld not track down. The error only seems to occur on GPU, not 
on CPU. When type=float, everything is fine as well.

I am using CUDA 4.0 on a GTX480 on gcc/Ubuntu.

Original issue reported on code.google.com by boiling....@gmail.com on 4 Jul 2011 at 1:23

[GoogleCodeExporter]

This is the expected behavior.

unsigned char cannot represent values greater than 255.

Original comment by jaredhoberock on 5 Jul 2011 at 8:51

• Changed state: Invalid

[GoogleCodeExporter]

The operator returns ints, though. thrust::reduce should return a value which 
is the same as the return type of the operator. On CPU it actually does that. 
Therefore, I don't quite understand your reasoning, can you explain in more 
detail?

Original comment by boiling....@gmail.com on 8 Sep 2011 at 8:45

[GoogleCodeExporter]

Unlike std::accumulate, thrust::reduce requires binary_op to be commutative 
[1].  Your expectation

  sum=binary_op(sum,nextval),

ignores commutativity.  The CUDA implementation is putting the sum on the RHS 
of the operator [2] and converting to unsigned char.

[1] 
http://wiki.thrust.googlecode.com/hg/html/group__reductions.html#ga268b8b31c9c3a
2e6e4b3b6be1f5c202d
[2] 
http://code.google.com/p/thrust/source/browse/thrust/detail/backend/cuda/reduce.
inl#99

Original comment by jaredhoberock on 8 Sep 2011 at 9:08

[GoogleCodeExporter]

Ok, I missed that point. Thanks for the explanation!

Original comment by boiling....@gmail.com on 8 Sep 2011 at 12:19