thrust::min_element doesn't work on a float4 device_vector, while it does on a host_vector

Question

I'm trying to find the minimum number in a array using Thrust and CUDA.
The following device example returns with 0 :

thrust::device_vector<float4>::iterator it =  thrust::min_element(IntsOnDev.begin(),IntsOnDev.end(),equalOperator());       
int pos = it - IntsOnDev.begin();

However, this host version works perfectly:

thrust::host_vector<float4>arr = IntsOnDev;
thrust::host_vector<float4>::iterator it2 =  thrust::min_element(arr.begin(),arr.end(),equalOperator());
int pos2 = it2 - arr.begin();

the comperator type :

struct equalOperator
{
  __host__ __device__
    bool operator()(const float4 x,const float4 y) const
    {
        return ( x.w < y.w );
    }
};

I just wanted to add that thrust::sort works with the same predicate.

Answer 1

Unfortunately, nvcc disagrees with some host compilers (some 64 bit versions of MSVC, if I recall correctly) about the size of certain aligned types. float4 is one of these. This often results in undefined behavior.

The work-around is to use types without alignment, for example my_float4:

struct my_float4
{
  float x, y, z, w;
};