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.

What happens if you try this with my_float4, i.e. struct my_float4 { float x,y,z,w; }; ? — Jared Hoberock
that did the trick! i found out that solution few mins after i wrote the original post... i actually defined a new float4 struct of my own, sine the new struct is basically equal to float4(in bytes) so i had to change almost nothing — igal k

Jared Hoberock Jared Hoberock · Accepted Answer · 2012-02-13T19:19:19

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;
};

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

1 Answers