CUDA functor of min point

Question

Im stunned. I readed a lot of answers and still can't get it work. Im trying to implement comparator over struct point, which will help to find minimal point in array, this is code so far:

struct minPointOperator
{
    __device__ __host__ point operator()(const point& x, const point& y) const 
    { 
        return x.value > y.value ? y : x;
    } 
};
int findBestPointIndx(point* dev_points, int pointCount)
{
    thrust::device_ptr<point> points(dev_points);
    int id = thrust::reduce(points, points+pointCount, 0, minPointOperator());
    return id;
}

but thi doesn't compile, it just spew a lot of function "minPointOperator::operator()" cannot be called with the given argument list errors

Answer 1

It may be helpful if you study the thrust documentation for reduce.

Note the following:

Template Parameters

InputIterator is a model of Input Iterator and InputIterator's value_type is convertible to T.

T is a model of Assignable, and is convertible to BinaryFunction's first_argument_type and second_argument_type.

This means that your choice of the init parameter as well as the return type of reduce are not arbitrary. They must be the same type (effectively) as the type of the input iterator (i.e. point in this case).

thrust::reduce (in this case) does not return the index of the minimum element. It returns the actual contents of the minimum element, i.e. if your elements are point, it returns a point, and the init value must be of type point.

To find the index of the minimum point (appears to be your objective) there are various possibilities. One possibility would be to zip together (using zip_iterator) the data along with an index array (e.g. counting_iterator), do the min-reduction based on that, and then extract the index from the value returned by thrust::reduce.

However, I think a better approach would be what is suggested here, i.e. use thrust::min_element which will return an interator to the minimum element, which can then be used to compute the index easily.

Here's a worked example:

$ cat t683.cu
#include <thrust/device_ptr.h>
#include <thrust/extrema.h>
#include <thrust/host_vector.h>
#include <thrust/device_vector.h>

#include <iostream>

struct point{

  int value;
  int data;
};


struct minPointOperator
{
    __device__ __host__ bool operator()(const point& x, const point& y) const
    {
        return (x.value < y.value);
    }
};

int findBestPointIndx(point* dev_points, int pointCount)
{
    thrust::device_ptr<point> points(dev_points);
    int id = thrust::min_element(points, points+pointCount, minPointOperator()) - points;
    return id;
}

int main(){

  point data[4] = { {4,5}, {0,7}, {2,3}, {6,1} };
  thrust::host_vector<point> h_data(data, data + 4);
  thrust::device_vector<point> d_data = h_data;
  int min_point = findBestPointIndx(thrust::raw_pointer_cast(d_data.data()), 4);
  std::cout << "min index: " << min_point << std::endl;
  return 0;
}

$ nvcc t683.cu -o t683
$ ./t683
min index: 1
$