CUDA: Getting max value and its index in an array

Question

I have several blocks were each block executes on separate part of an integer array. As an example: block one from array[0] to array[9] and block two from array[10] to array[20].

What is the best way i can get the index of the max value of the array for each block?

Example block one a[0] to a[10] have the following values:
5 10 2 3 4 34 56 3 9 10

So 56 is the largest value at index 6.

I cannot use the shared memory because the size of the array may be very big. Therefore it won't fit. Are there any libraries that allows me to do so fast?

I know about the reduction algorithm, but i think my case is different because i want to get the index of the largest element.

Answer 1

If I understood exactly what you want is : Get the index for the array A of the max value inside it.

If that is true then I would suggest you to use the thrust library:

Here is how you would do it:

#include <thrust/device_vector.h>
#include <thrust/tuple.h>
#include <thrust/reduce.h>
#include <thrust/fill.h>
#include <thrust/generate.h>
#include <thrust/sort.h>
#include <thrust/sequence.h>
#include <thrust/copy.h>
#include <cstdlib>
#include <time.h>

using namespace thrust;

// return the biggest of two tuples
template <class T>
struct bigger_tuple {
    __device__ __host__
    tuple<T,int> operator()(const tuple<T,int> &a, const tuple<T,int> &b) 
    {
        if (a > b) return a;
        else return b;
    } 

};

template <class T>
int max_index(device_vector<T>& vec) {

    // create implicit index sequence [0, 1, 2, ... )
    counting_iterator<int> begin(0); counting_iterator<int> end(vec.size());
    tuple<T,int> init(vec[0],0); 
    tuple<T,int> smallest;

    smallest = reduce(make_zip_iterator(make_tuple(vec.begin(), begin)), make_zip_iterator(make_tuple(vec.end(), end)),
                      init, bigger_tuple<T>());
    return get<1>(smallest);
}

int main(){

    thrust::host_vector<int> h_vec(1024);
    thrust::sequence(h_vec.begin(), h_vec.end()); // values = indices

    // transfer data to the device
    thrust::device_vector<int> d_vec = h_vec;

    int index = max_index(d_vec);

    std::cout <<  "Max index is:" << index <<std::endl;
    std::cout << "Value is: " << h_vec[index] <<std::endl;

    return 0;
}

Answer 2

This will not benefit the original poster but for those who came to this page looking for an answer I would second the recommendation to use thrust that already has a function thrust::max_element that does exactly that - returns an index of the largest element. min_element and minmax_element functions are also provided. See thrust documentation for details here.

Answer 3

As well as the suggestion to use Thrust, you could also use the CUBLAS cublasIsamax function.