CUDA: получение максимального значения и его индекса в массиве

Question

У меня есть несколько блоков, каждый из которых выполняется в отдельной части целочисленного массива. В качестве примера: блокируйте один из массива [0] в массив [9] и блокируйте два из массива [10] в массив [20].

Как лучше всего получить индекс максимального значения массива для каждого блока?

Пример первого блока a [0] - a [10] имеет следующие значения:
5 10 2 3 4 34 56 3 9 10

То есть 56 - самое большое значение в индексе 6.

Я не могу использовать разделяемую память, потому что размер массива может быть очень большим. Поэтому оно не подходит. Есть ли библиотеки, которые позволяют мне делать это так быстро?

Я знаю об алгоритме редукции, но я думаю, что мой случай другой, потому что я хочу получить индекс наибольшего элемента.

Answer 1

Если я точно понял, что вы хотите: получить индекс для массива A с максимальным значением внутри него.

Если это так, то я бы предложил вам использовать библиотеку тяги:

Вот как бы вы это сделали:

#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

Это не пойдет на пользу оригинальному постеру, но для тех, кто пришел на эту страницу в поисках ответа, я бы рекомендовал использовать thrust, который уже имеет функцию thrust :: max_element, которая делает именно это - возвращает индекс наибольшегоэлемент.Функции min_element и minmax_element также предоставляются.Подробнее см. Документацию по тяге здесь .

Answer 3

Помимо предложения использовать Thrust, вы также можете использовать функцию CUBLAS cublasIsamax.

Answer 4

Одна вещь, на которую следует обратить внимание при выполнении максимального значения плюс уменьшение индекса, это то, что если в вашем массиве более одного идентичного максимального элемента, т.е. в вашем примере, если было 2 или более значений, равных 56,Возвращаемый индекс не будет уникальным и, возможно, будет отличаться при каждом запуске кода, потому что время упорядочения потока в графическом процессоре не является детерминированным.

Чтобы обойти эту проблему, вы можете использовать уникальный индекс порядка, такой как threadid + threadsperblock * blockid, или же расположение индекса элемента, если он уникален.Тогда максимальный тест выполняется по следующим направлениям:

if(a>max_so_far || a==max_so_far && order_a>order_max_so_far)
{ 
    max_so_far = a;
    index_max_so_far = index_a;
    order_max_so_far = order_a;
}

(индекс и порядок могут быть одной и той же переменной, в зависимости от приложения.)

Answer 5

Размер вашего массива по сравнению с разделяемой памятью практически не имеет значения, поскольку количество потоков в каждом блоке является ограничивающим фактором, а не размером массива. Одно из решений состоит в том, чтобы каждый блок потока работал с размером массива, равным размеру блока потока. То есть, если у вас 512 потоков, тогда блок n будет искать массив [n] - массив [n + 511]. Каждый блок делает сокращение, чтобы найти самый высокий член в этой части массива. Затем вы возвращаете максимум каждого раздела обратно на хост и выполняете простой линейный поиск, чтобы найти наибольшее значение в общем массиве. Каждое сокращение без GPU уменьшает линейный поиск в 512 раз. В зависимости от размера массива, вы можете захотеть сделать больше сокращений, прежде чем возвращать данные. (Если ваш массив имеет размер 3 * 512 ^ 10, вы можете сделать 10 сокращений в gpu и выполнить поиск хоста через 3 оставшиеся точки данных.)