As I mentioned in my comment above, I would recommend to avoid writing reductions of your own and use CUDA Thrust whenever possible. This holds true even in the case when you need to customize those operations, the customization being possible by properly overloading, e.g., relational operations.
Below I'm providing a simple code to evaluate the minimum in an array along with its index. It is based on a classical example contained in the An Introduction to Thrust presentation. The only addition is skipping, as you requested, the -1
's from the counting. This can be reasonably done by replacing all the -1
's in the array by INT_MAX
, i.e., the maximum representable integer according to IEEE floating point standards.
#include <thrust\device_vector.h>
#include <thrust\replace.h>
#include <thrust\sequence.h>
#include <thrust\reduce.h>
#include <thrust\iterator\zip_iterator.h>
#include <thrust\tuple.h>
// --- Struct returning the smallest of two tuples
struct smaller_tuple
{
__host__ __device__ thrust::tuple<int,int> operator()(thrust::tuple<int,int> a, thrust::tuple<int,int> b)
{
if (a < b)
return a;
else
return b;
}
};
void main() {
const int N = 20;
const int large_value = INT_MAX;
// --- Setting the data vector
thrust::device_vector<int> d_vec(N,10);
d_vec[3] = -1; d_vec[5] = -2;
// --- Copying the data vector to a new vector where the -1's are changed to FLT_MAX
thrust::device_vector<int> d_vec_temp(d_vec);
thrust::replace(d_vec_temp.begin(), d_vec_temp.end(), -1, large_value);
// --- Creating the index sequence [0, 1, 2, ... )
thrust::device_vector<int> indices(d_vec_temp.size());
thrust::sequence(indices.begin(), indices.end());
// --- Setting the initial value of the search
thrust::tuple<int,int> init(d_vec_temp[0],0);
thrust::tuple<int,int> smallest;
smallest = thrust::reduce(thrust::make_zip_iterator(thrust::make_tuple(d_vec_temp.begin(), indices.begin())),
thrust::make_zip_iterator(thrust::make_tuple(d_vec_temp.end(), indices.end())),
init, smaller_tuple());
printf("Smallest %i %i\n",thrust::get<0>(smallest),thrust::get<1>(smallest));
getchar();
}